Benchmark · pre-generated, always available← back to leaderboard

Migration Readiness Report

Dao-AILab/flash-attention

https://github.com/Dao-AILab/flash-attention ↗

68/ 100 ready

Manual blockers present

Advisory Summary — Dao-AILab/flash-attention → ROCm Migration

With a readiness score of 68/100, this repo is roughly two-thirds of the way to ROCm compatibility: 731 findings already work as-is and 1,076 are mechanical replacements (e.g., cuda→hip API renames, header paths, built-in thread-index substitutions) that can be batch-processed with low risk. The real blocker is the 47 manual findings — likely concentrated in custom kernel intrinsics (e.g., __shfl_xor_sync, warp-level primitives, cp.async / ldmatrix equivalents), TMA/descriptor usage, and build-system assumptions around nvcc/cutlass that require human judgment and architecture-specific rewrites for RDNA/CDNA. First step: triage those 47 manual blockers by file to identify which kernels depend on features with no direct HIP equivalent, then scope whether to rewrite those paths using ROCm-native primitives (e.g., __shfl_xor, buffer_load, MFMA intrinsics) or to gate them behind conditional compilation while the mechanical sweep proceeds in parallel.

731Works as-is

1076Mechanical change

47Manual blocker

Findings

1854

Python files

222

Files scanned

1072

Custom CUDA kernels

detected

Detected but out of scope (not analyzed): C++, C/C++ header

Findings by file

1854 findings · 759 files

AI/racecheck_repro_1d_bulk.py· 4

cuda-python low-level driver API

import cuda.bindings.driver as cuda

Fireworks AI:16

The cuda.bindings.driver module provides low-level CUDA driver API access in Python; its ROCm counterpart is amdsmi for device/system management or rocprofiler-sdk for profiling-related driver interactions. Given the racecheck context, rocprofiler-sdk is the most relevant replacement, but the API surface is not 1:1 and call sites will need individual review.

RecommendRewrite driver-API calls against the HIP runtime.

Suggested change · advisory

--- AI/racecheck_repro_1d_bulk.py
+++ AI/racecheck_repro_1d_bulk.py
@@ -16,1 +16,1 @@
-import cuda.bindings.driver as cuda
+# Advisory: replace with rocprofiler-sdk for race-check/profiling, or amdsmi for device management
+import rocprofiler_sdk as rocm  # was: import cuda.bindings.driver as cuda

Device string "cuda"

src = torch.arange(TILE * N_BLKS, device="cuda", dtype=torch.float32)

template:68

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

cuda.CUstream(torch.cuda.current_stream().cuda_stream))

template:71

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:72

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

AI/racecheck_repro_1d_tensor.py· 4

cuda-python low-level driver API

import cuda.bindings.driver as cuda

Fireworks AI:16

This import uses cuda-python's low-level driver bindings, which have no direct ROCm equivalent in the same package namespace. Migration requires switching to the ROCm Python bindings (e.g., amd.hip from rocm-python) and translating driver API calls to HIP equivalents, which often differ in naming and handle types.

RecommendRewrite driver-API calls against the HIP runtime.

Suggested change · advisory

--- AI/racecheck_repro_1d_tensor.py
+++ AI/racecheck_repro_1d_tensor.py
@@ -16,1 +16,1 @@
-import cuda.bindings.driver as cuda
+import amd.hip as hip  # advisory: replace cuda-python driver API with rocm-python HIP bindings

Device string "cuda"

src = torch.arange(TILE * N_BLKS, device="cuda", dtype=torch.float32)

template:78

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

cuda.CUstream(torch.cuda.current_stream().cuda_stream))

template:81

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:82

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/cute/cute_dsl_ptxas.py· 15

Inline PTX assembly

System ptxas replacement for CUTLASS DSL.

Fireworks AI:2

This file wraps/replaces the system ptxas binary, which is NVIDIA-specific and has no ROCm equivalent; the ROCm compilation path uses amdclang/clang and emits GCN ISA rather than PTX. Any CUTLASS DSL logic that invokes ptxas or parses PTX must be redirected to the ROCm compiler toolchain or stubbed out. Inline PTX/ptxas dependencies will not function on AMD Instinct GPUs.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
@@ -2,3 +2,8 @@
-# System ptxas replacement for CUTLASS DSL.
+# Advisory: ptxas is NVIDIA-only; on ROCm there is no ptxas equivalent.
+# The CUTLASS DSL compilation path must be adapted to use the ROCm
+# compiler toolchain (amdclang/clang) emitting GCN ISA, or this
+# wrapper must be stubbed/disabled when targeting AMD Instinct GPUs.
+# Consider detecting the backend (CUDA vs HIP) and branching accordingly.

Inline PTX assembly

CUTE_DSL_PTXAS_PATH - Path to ptxas (e.g., /usr/local/cuda/bin/ptxas)

Fireworks AI:4

This file references ptxas, NVIDIA's PTX assembler, which does not exist in the ROCm toolchain. Any inline PTX assembly validated or compiled through this path must be replaced with ROCm-equivalent intrinsics or GCN ISA, and the ptxas path configuration becomes invalid on AMD Instinct GPUs.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
-CUTE_DSL_PTXAS_PATH    - Path to ptxas (e.g., /usr/local/cuda/bin/ptxas)
+CUTE_DSL_PTXAS_PATH    - Advisory: ptxas is NVIDIA-only; ROCm has no equivalent.
+                       Inline PTX must be ported to HIP intrinsics or GCN ISA.

Inline PTX assembly

print(f"[ptxas] {msg}", file=sys.stderr)

Fireworks AI:27

The snippet itself is only a log/print line referencing ptxas (NVIDIA's PTX assembler) and contains no inline PTX, so there is no direct assembly to translate. However, if this file shells out to ptxas elsewhere, that invocation must be replaced with ROCm's assembler path (e.g., hipcc/clang with -x assembler or comgr), since ptxas does not exist on AMD Instinct. The log string is cosmetic but should be generalized to avoid implying an NVIDIA-only toolchain.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
@@ -24,3 +24,3 @@
-    print(f"[ptxas] {msg}", file=sys.stderr)
+    print(f"[asm] {msg}", file=sys.stderr)

Inline PTX assembly

for ptx_path in Path(dump_dir).glob(f"*{func_name}*.ptx"):

Fireworks AI:37

PTX is NVIDIA-specific and has no direct ROCm equivalent; ROCm uses GCN/CDNA ISA. This code globs for .ptx files emitted by ptxas, which will never exist on AMD toolchains. The migration requires replacing PTX parsing with GCN ISA disassembly (e.g., via rocobjdump or amdgcn-isa output) and adjusting downstream parsing accordingly.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
@@ -37,1 +37,1 @@
-for ptx_path in Path(dump_dir).glob(f"*{func_name}*.ptx"):
+# Advisory: ROCm does not emit PTX; consider globbing for GCN ISA files
+# (e.g. "*.s" or "*.isa") produced by rocobjdump / --save-temps instead.
+for ptx_path in Path(dump_dir).glob(f"*{func_name}*.s"):

Inline PTX assembly

"""Compile PTX to cubin using system ptxas."""

Fireworks AI:46

PTX is NVIDIA-specific and cannot be compiled or executed on AMD Instinct GPUs; ptxas has no ROCm equivalent. This functionality would need to be replaced with GCN/ROCm ISA assembly paths or removed in favor of HIP-native code generation. Migration impact is high since any downstream PTX kernels will fail on ROCm.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

- """Compile PTX to cubin using system ptxas."""
+ """Compile PTX to cubin using system ptxas."""  # ADVISORY: PTX/ptxas are NVIDIA-only; replace with ROCm GCN ISA path or HIP codegen for AMD Instinct migration.

Inline PTX assembly

raise RuntimeError(f"ptxas failed: {result.stderr}")

Fireworks AI:65

This file invokes ptxas, the NVIDIA PTX assembler, which does not exist in the ROCm toolchain. On AMD Instinct GPUs, inline PTX must be replaced with GCN/AMDGPU ISA or HIP intrinsics, and any ptxas-based validation step must be removed, stubbed, or routed to an ROCm-aware path. This is advisory only and requires manual porting of the surrounding PTX generation logic.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
@@ -62,6 +62,12 @@
 def _run_ptxas(ptx_path, arch, output_path):
+    # ADVISORY: ptxas is NVIDIA-only and unavailable on ROCm.
+    # On AMD Instinct, inline PTX must be ported to GCN/AMDGPU ISA
+    # or HIP intrinsics. Consider gating this path behind a
+    # backend check (e.g., torch.version.hip) and providing an
+    # ROCm-aware validation path or skipping validation.
     result = subprocess.run(["ptxas", ...], capture_output=True)
     if result.returncode != 0:
-        raise RuntimeError(f"ptxas failed: {result.stderr}")
+        raise RuntimeError(f"ptxas failed (NVIDIA-only; not available on ROCm): {result.stderr}")

Inline PTX assembly

"""Replacement for _load_cuda_library that uses system ptxas."""

Fireworks AI:82

This helper loads NVIDIA's ptxas, the PTX-to-SASS assembler, which has no ROCm equivalent; AMD toolchains use amd_comgr/clang to assemble GCN/CDNA ISA. Any downstream code that feeds inline PTX through this path will not function on Instinct GPUs and must be replaced with HIP intrinsics or GCN inline asm. Advisory only — no automatic fix is possible without rewriting the PTX fragments themselves.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- a/flash_attn/cute/cute_dsl_ptxas.py
+++ b/flash_attn/cute/cute_dsl_ptxas.py
@@ -82,3 +82,8 @@
 """Replacement for _load_cuda_library that uses system ptxas."""
+# Advisory: ptxas is NVIDIA-only. On ROCm, inline PTX cannot be assembled;
+# replace PTX fragments with HIP intrinsics or GCN inline asm, and route
+# assembly through amd_comgr / clang instead of this loader.
+import os
+_IS_ROCM = os.environ.get('HIP_PLATFORM', '') == 'amd' or os.path.exists('/opt/rocm')
+if _IS_ROCM:
+    raise RuntimeError('ptxas loader is CUDA-only; use amd_comgr on ROCm')

Inline PTX assembly

_log("PTX not found, falling back to embedded ptxas")

Fireworks AI:86

PTX and ptxas are NVIDIA-specific; AMD Instinct GPUs use GCN/CDNA ISA and do not accept PTX input. Any embedded ptxas fallback path will fail on ROCm and must be replaced with an AMDGPU assembler path or removed. This is advisory only—manual validation of the replacement strategy is required.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
@@ -86,1 +86,1 @@
-_log("PTX not found, falling back to embedded ptxas")
+_log("ROCm: PTX/ptxas unavailable; AMDGPU ISA path required (advisory)")

Inline PTX assembly

_log(f"Compilation failed ({e}), falling back to embedded ptxas")

Fireworks AI:94

This file relies on ptxas (NVIDIA's PTX assembler) as a compilation fallback, which does not exist in the ROCm toolchain. On AMD Instinct GPUs, PTX is unavailable; the equivalent low-level path would use AMDGPU ISA via clang/lld or the ROCm offline compiler (amdgcn). Any embedded PTX blobs or ptxas invocation logic in this module must be replaced or conditionally bypassed for ROCm builds.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
@@ -91,6 +91,8 @@
             except Exception as e:
-                _log(f"Compilation failed ({e}), falling back to embedded ptxas")
+                if not _IS_ROCM:
+                    _log(f"Compilation failed ({e}), falling back to embedded ptxas")
+                else:
+                    raise RuntimeError(f"Compilation failed ({e}); ptxas fallback unsupported on ROCm") from e

cuda-python low-level driver API

import cuda.bindings.runtime as cuda_runtime

Fireworks AI:98

The cuda-python bindings (cuda.bindings.runtime) provide Python-level access to the CUDA runtime/driver API and have no direct ROCm equivalent. This module is likely used for PTX/JIT compilation orchestration in the cute DSL layer; on ROCm, the equivalent functionality would go through HIP compilation APIs (e.g., hiprtc) or shell-out to clang/amdgcn tooling, but no drop-in Python binding package exists today.

RecommendRewrite driver-API calls against the HIP runtime.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
@@ -95,7 +95,11 @@
 # Advisory: cuda-python bindings are CUDA-only; guard or replace with HIP/ROCm
 # compilation tooling (hiprtc or clang/amdgcn) when targeting Instinct GPUs.
- import cuda.bindings.runtime as cuda_runtime
+ try:
+     import cuda.bindings.runtime as cuda_runtime
+ except ImportError:
+     cuda_runtime = None  # ROCm path: use hiprtc or external compiler invocation

Inline PTX assembly

_log(f"cudaLibraryLoadData failed ({err}), falling back to embedded ptxas")

Fireworks AI:102

This log message references ptxas, NVIDIA's PTX-to-SASS assembler, which has no ROCm equivalent. The surrounding code that invokes ptxas or uses cudaLibraryLoadData (CUDA 12+ lazy module loading) must be replaced with ROCm's compilation pipeline (e.g., hipModuleLoadData or offline compilation via clang/llvm to GCN ISA). Any embedded PTX assembly strings elsewhere in this file are NVIDIA-specific and must be rewritten as AMDGPU GCN assembly or replaced with HIP intrinsics.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

- _log(f"cudaLibraryLoadData failed ({err}), falling back to embedded ptxas")
+ _log(f"hipModuleLoadData failed ({err}), falling back to embedded amdgcn compilation")

Inline PTX assembly

_log("cuda_load_to_device failed, falling back to embedded ptxas")

Fireworks AI:120

This file invokes ptxas, the NVIDIA PTX assembler, which has no ROCm equivalent; AMD GPUs use GCN/AMDGPU ISA assembled via clang/llvm rather than PTX. Any PTX generation or ptxas fallback path must be replaced with ROCm-native assembly tooling or HIP intrinsics, or the DSL layer must be ported to emit AMDGPU ISA. The log line itself is benign, but it signals a code path that will fail on Instinct GPUs.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Inline PTX assembly

"""Install system ptxas hook. Call before importing cutlass."""

Fireworks AI:133

This file manages ptxas, the NVIDIA PTX assembler, which has no ROCm equivalent. On AMD Instinct GPUs, PTX assembly is not supported; the equivalent low-level ISA is GCN/RDNA, assembled via the ROCm toolchain. Any code paths that invoke ptxas or embed PTX must be stubbed, replaced with HIP/ROCm equivalents, or guarded behind a backend check.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- a/flash_attn/cute/cute_dsl_ptxas.py
+++ b/flash_attn/cute/cute_dsl_ptxas.py
@@ -130,7 +130,11 @@
 """Install system ptxas hook. Call before importing cutlass."""
+# Advisory: ptxas is NVIDIA-only and has no ROCm equivalent.
+# On AMD Instinct, this hook should be a no-op or replaced with
+# a ROCm assembler hook if inline GCN assembly is needed.
+if os.environ.get('HIP_PLATFORM') == 'amd':
+    return  # No-op on ROCm; ptxas does not exist.

Inline PTX assembly

raise RuntimeError(f"ptxas not found: {CUTE_DSL_PTXAS_PATH}")

Fireworks AI:138

This code path hard-depends on NVIDIA's ptxas (the PTX assembler shipped only with the CUDA Toolkit), which does not exist on ROCm. On AMD Instinct, any CUTE DSL flow that invokes ptxas for PTX validation/compilation will unconditionally raise this RuntimeError and must be replaced with an AMDGPU-aware path (e.g., clang/lld targeting amdgcn) or bypassed entirely.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
@@ -135,6 +135,12 @@
     if not os.path.isfile(CUTE_DSL_PTXAS_PATH):
-        raise RuntimeError(f"ptxas not found: {CUTE_DSL_PTXAS_PATH}")
+        # Advisory: on ROCm there is no ptxas; fall back to clang/lld for AMDGPU
+        # or skip PTX validation entirely. Replace with an ROCm-aware assembler
+        # path before enabling this flow on AMD Instinct.
+        if os.environ.get("HIP_PLATFORM") or os.environ.get("ROCM_PATH"):
+            CUTE_DSL_PTXAS_PATH = None  # disable ptxas-dependent validation
+        else:
+            raise RuntimeError(f"ptxas not found: {CUTE_DSL_PTXAS_PATH}")

Inline PTX assembly

"Require CUTE_DSL_KEEP_PTX=1 to use system's ptxas"

Fireworks AI:144

This code references ptxas, the NVIDIA PTX assembler, which does not exist in the ROCm toolchain. On AMD Instinct GPUs, PTX intermediate assembly has no equivalent; ROCm uses LLVM IR and GCN/CDNA ISA, so any CUTE DSL pipeline that invokes ptxas must be redirected to the ROCm assembler (llvm-mc or clang with amdgcn target) or the PTX generation path must be bypassed entirely.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_ptxas.py
+++ flash_attn/cute/cute_dsl_ptxas.py
@@ -142,3 +142,8 @@
-"Require CUTE_DSL_KEEP_PTX=1 to use system's ptxas"
+"Require CUTE_DSL_KEEP_PTX=1 to use system's ptxas"
+# Advisory: ptxas is NVIDIA-only. On ROCm, replace ptxas invocation with
+# `clang -target amdgcn-amd-amdhsa` or `llvm-mc -triple=amdgcn-amd-amdhsa`.
+# PTX assembly itself has no ROCm equivalent; the CUTE DSL PTX path must be
+# stubbed or ported to GCN inline asm / LLVM IR.

flash_attn/cute/cute_dsl_utils.py· 4

Triton dependency

from triton.tools.disasm import extract

template:9

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

torch.cuda API usage

return torch.cuda.get_device_capability(device)

template:41

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

cuda-python low-level driver API

from cuda.bindings import driver

Fireworks AI:128

The cuda.bindings.driver import provides low-level CUDA driver API access (device/context/module management) with no drop-in ROCm Python equivalent. Migration requires either using HIP C/C++ via ctypes bindings, pyrsmi/rocm_smi for device queries, or rewriting the affected logic in HIP and exposing it through a Python extension. This is a manual porting task and cannot be auto-translated.

RecommendRewrite driver-API calls against the HIP runtime.

Suggested change · advisory

--- flash_attn/cute/cute_dsl_utils.py
+++ flash_attn/cute/cute_dsl_utils.py
@@ -128,1 +128,1 @@
-from cuda.bindings import driver
+# ROCm migration advisory: no direct Python equivalent for cuda.bindings.driver.
+# Replace with HIP/ROCm bindings (e.g., ctypes to HIP runtime, or rocm_smi for
+# device queries). Manual porting required.

torch.cuda API usage

device_id = torch.cuda.current_device()

template:132

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/cute/flash_bwd_mla_dk_sm100.py· 2

cuda-python low-level driver API

import cuda.bindings.driver as cuda

Fireworks AI:31

The cuda.bindings.driver module is part of cuda-python and provides low-level CUDA driver API access (context, module, memory, stream management). ROCm has no drop-in Python equivalent with the same API surface; the closest analog is amdsmi for device/system queries or HIP runtime bindings via PyTorch's HIP backend. This import and all downstream driver calls must be replaced or stubbed with ROCm-compatible equivalents, and the surrounding SM100-specific CUTLUTE code is NVIDIA-architecture-specific and unlikely to port directly to AMD Instinct GPUs.

RecommendRewrite driver-API calls against the HIP runtime.

Suggested change · advisory

--- a/flash_attn/cute/flash_bwd_mla_dk_sm100.py
+++ b/flash_attn/cute/flash_bwd_mla_dk_sm100.py
@@ -28,7 +28,8 @@
 
 # Advisory: cuda.bindings.driver has no direct ROCm Python equivalent.
 # Replace with amdsmi for device queries or HIP runtime via PyTorch, and
 # audit all downstream cuda.* driver calls in this file.
-import cuda.bindings.driver as cuda
+# import cuda.bindings.driver as cuda
+import amdsmi as cuda  # placeholder — API surface differs; manual port required

FlashAttention dependency

from flash_attn.cute.utils import get_batch_from_cu_tensor

template:48

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_bwd_mla_dq_dqv_sm100.py· 3

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:37

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.topk_gather_kv import CpasyncGatherKVManager

template:47

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import get_batch_from_cu_tensor

template:48

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_bwd_mla_sm100.py· 11

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:7

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.pack_gqa import pack_gqa_layout

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_info import BlockInfo

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

import flash_attn.cute.blackwell_helpers as fa_sm100_utils

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.fa_logging import fa_log, fa_printf

template:33

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import smid, elem_pointer, get_batch_from_cu_tensor

template:34

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.copy_utils import tiled_copy_2d, atomic_add_fp32x4

template:35

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.topk_gather_kv import CpasyncGatherKVManager

template:37

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.named_barrier import NamedBarrierBwdSm100_MLA2CTA

template:40

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_bwd_postprocess.py· 6

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:7

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute import utils

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import assume_tensor_aligned

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import ampere_helpers as sm80_utils

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:27

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_bwd_preprocess.py· 5

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:19

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute import utils

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfo

template:29

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:31

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.pack_gqa import pack_gqa_layout

template:36

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_bwd_sm100.py· 16

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:6

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute import utils

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import assume_tensor_aligned

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import copy_utils

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import pipeline

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.blackwell_helpers import gemm_w_idx, gemm_ptx_w_idx # noqa

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import AttentionMask

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_info import BlockInfo

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import barrier

template:35

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.named_barrier import NamedBarrierBwdSm100

template:36

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.softmax import apply_score_mod_inner, apply_score_mod_bwd_inner

template:37

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import BlockSparseTensors

template:38

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:39

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparse_utils import (

template:40

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_bwd_sm90.py· 14

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:5

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import assume_tensor_aligned

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import utils

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import AttentionMask

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_info import BlockInfo

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import pipeline

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:27

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import barrier

template:33

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.named_barrier import NamedBarrierBwd

template:34

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.softmax import apply_score_mod_inner, apply_score_mod_bwd_inner

template:35

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import BlockSparseTensors

template:36

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:37

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparse_utils import (

template:38

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_bwd.py· 11

cuda-python low-level driver API

import cuda.bindings.driver as cuda

Fireworks AI:9

The cuda.bindings.driver import is the cuda-python low-level driver API, which has no drop-in ROCm equivalent. On AMD Instinct GPUs, the advisory migration is to amd.hip (hip-python) bindings, but the API surface differs and call sites must be reviewed individually. This is advisory only and requires manual validation of downstream usage.

RecommendRewrite driver-API calls against the HIP runtime.

Suggested change · advisory

--- a/flash_attn/cute/flash_bwd.py
+++ b/flash_attn/cute/flash_bwd.py
@@ -9,1 +9,1 @@
-import cuda.bindings.driver as cuda
+# Advisory: migrate to hip-python bindings; API surface differs, review call sites
+import amd.hip as hip  # noqa: F401  (advisory only, not auto-fixed)

FlashAttention dependency

from flash_attn.cute import ampere_helpers as sm80_utils

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import assume_tensor_aligned

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import utils

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import AttentionMask

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.softmax import call_score_mod, call_score_mod_bwd

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_info import BlockInfo

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import SingleTileScheduler, SingleTileVarlenScheduler, TileSchedulerArguments

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import BlockSparseTensors

template:27

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_fwd_combine.py· 4

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:8

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute import utils

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import assume_tensor_aligned

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfo

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_fwd_mla_sm100.py· 14

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:8

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.pack_gqa import pack_gqa_layout, make_packgqa_tiled_tma_atom

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.paged_kv import PagedKVManager

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import utils as fa_utils

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_info import BlockInfo

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import AttentionMask

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

import flash_attn.cute.blackwell_helpers as fa_sm100_utils

template:27

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.softmax import SoftmaxSm100

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:29

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.fa_logging import fa_log, fa_printf

template:39

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import smid, get_batch_from_cu_tensor

template:40

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.topk_gather_kv import CpasyncGatherKVManager

template:42

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.named_barrier import NamedBarrierFwdSm100_MLA2CTA

template:45

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_fwd_sm100.py· 19

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:20

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.paged_kv import PagedKVManager

template:36

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import assume_tensor_aligned

template:37

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import utils

template:38

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

import flash_attn.cute.pipeline as pipeline_custom

template:39

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import AttentionMask

template:41

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.softmax import SoftmaxSm100, apply_score_mod_inner

template:42

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:43

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_info import BlockInfo

template:44

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import BlockSparseTensors

template:45

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparse_utils import (

template:46

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.pack_gqa import PackGQA, pack_gqa_layout

template:52

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import mma_sm100_desc as sm100_desc

template:53

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import blackwell_helpers as sm100_utils

template:54

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.named_barrier import NamedBarrierFwdSm100

template:55

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:58

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.fa_logging import fa_log, fa_printf

template:68

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import smid

template:69

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:70

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_fwd_sm90.py· 16

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:8

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import assume_tensor_aligned

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import utils

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import AttentionMask

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.softmax import Softmax, apply_score_mod_inner

template:27

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_info import BlockInfo

template:29

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import BlockSparseTensors

template:30

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparse_utils import (

template:31

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import pipeline as pipeline_custom

template:35

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.pack_gqa import PackGQA, pack_gqa_layout, make_packgqa_tiled_tma_atom

template:36

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.paged_kv import PagedKVManager

template:37

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.named_barrier import NamedBarrierFwd

template:38

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:40

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd import FlashAttentionForwardBase

template:48

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:49

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_fwd.py· 14

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:13

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute import ampere_helpers as sm80_utils

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import assume_tensor_aligned

template:27

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import utils

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import AttentionMask

template:29

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.softmax import Softmax, apply_score_mod_inner

template:30

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:31

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_info import BlockInfo

template:32

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.pack_gqa import PackGQA, pack_gqa_layout

template:33

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.named_barrier import NamedBarrierFwd

template:34

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import BlockSparseTensors

template:35

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import SingleTileScheduler, SingleTileVarlenScheduler, TileSchedulerArguments

template:36

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:37

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd_sm90 import FlashAttentionForwardSm90

template:1241

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/interface.py· 28

FlashAttention dependency

from flash_attn.cute.cache_utils import get_jit_cache

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.testing import is_fake_mode

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import cute_dsl_ptxas # noqa: F401

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Inline PTX assembly

# Patch to dump ptx and then use system ptxas to compile to cubin

template:25

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

FlashAttention dependency

from flash_attn.cute import utils

template:29

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import fa_logging

template:30

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import (

template:31

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd import FlashAttentionForwardSm80

template:37

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd_sm90 import FlashAttentionForwardSm90

template:38

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd_sm100 import FlashAttentionForwardSm100, DescaleTensors

template:39

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd_sm120 import FlashAttentionForwardSm120

template:40

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_bwd_preprocess import FlashAttentionBackwardPreprocess

template:41

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_bwd import FlashAttentionBackwardSm80

template:42

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_bwd_sm90 import FlashAttentionBackwardSm90

template:43

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_bwd_sm100 import FlashAttentionBackwardSm100

template:44

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_bwd_sm120 import FlashAttentionBackwardSm120

template:45

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_bwd_postprocess import FlashAttentionBackwardPostprocess

template:46

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd_combine import FlashAttentionForwardCombine

template:47

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd_mla_sm100 import FlashAttentionMLAForwardSm100

template:48

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_bwd_mla_sm100 import FlashAttentionSparseMLABackwardSm100

template:49

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_bwd_mla_dq_dqv_sm100 import dQdQvGemmKernel

template:50

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_bwd_mla_dk_sm100 import dKGemmKernel

template:51

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.sm100_hd256_2cta_fmha_forward import BlackwellFusedMultiHeadAttentionForward

template:54

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.sm100_hd256_2cta_fmha_backward import BlackwellFusedMultiHeadAttentionBackward

template:55

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:57

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import (

template:58

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

major, minor = torch.cuda.get_device_capability()

template:91

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

num_SMs = 132 if is_fake_mode() else torch.cuda.get_device_properties(device).multi_processor_count

template:566

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/cute/sm100_hd256_2cta_fmha_backward_dkdvkernel.py· 3

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:14

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

import flash_attn.cute.copy_utils as fa_copy_utils

template:35

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/sm100_hd256_2cta_fmha_backward_dqkernel.py· 5

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:5

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import (

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import SM100_TMEM_CAPACITY_COLUMNS

template:31

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

import flash_attn.cute.copy_utils as fa_copy_utils

template:32

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/sm100_hd256_2cta_fmha_backward.py· 5

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:12

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.sm100_hd256_2cta_fmha_backward_dqkernel import (

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.sm100_hd256_2cta_fmha_backward_dkdvkernel import (

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import assume_tensor_aligned

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData, as_bshkrd_tensor, as_shhb_tensor

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/sm100_hd256_2cta_fmha_forward.py· 6

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:6

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import (

template:27

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import SM100_TMEM_CAPACITY_COLUMNS

template:30

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd_sm100 import DescaleTensors, _TUNING_CONFIG

template:31

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import ex2_emulation_2, as_bshkrd_tensor, AuxData

template:32

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

hopper/setup.py· 9

Device string "cuda"

if BUILD_TARGET == "cuda":

template:97

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Inline PTX assembly

NVIDIA_TOOLCHAIN_VERSION = {"nvcc": "12.6.85", "ptxas": "12.8.93"}

template:438

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Inline PTX assembly

# ptxas 12.8 gives the best perf currently

template:467

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Inline PTX assembly

name="ptxas",

template:481

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Inline PTX assembly

src_func=lambda system, arch, version: f"cuda_nvcc-{system}-{arch}-{version}-archive/bin/ptxas",

template:482

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Inline PTX assembly

version=NVIDIA_TOOLCHAIN_VERSION["ptxas"],

template:484

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Inline PTX assembly

name="ptxas",

template:489

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Inline PTX assembly

version=NVIDIA_TOOLCHAIN_VERSION["ptxas"],

template:492

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Inline PTX assembly

# "--ptxas-options=--verbose,--register-usage-level=5,--warn-on-local-memory-usage", # printing out number of registers

template:620

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

setup.py· 5

Device string "cuda"

if BUILD_TARGET == "cuda":

template:50

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Inline PTX assembly

# "--ptxas-options=-v",

template:322

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

Inline PTX assembly

# "--ptxas-options=-O2",

template:323

Inline PTX is NVIDIA ISA and cannot run on AMD. The block must be hand-rewritten in HIP/GCN or replaced with a portable path.

RecommendRewrite the PTX block in HIP or a portable high-level equivalent.

torch.cuda API usage

if not torch.cuda.is_available():

template:472

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

props = torch.cuda.get_device_properties(torch.cuda.current_device())

template:477

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/cute/benchmark_mask_mod.py· 13

cuda-python low-level driver API

import cuda.bindings.driver as cuda

template:10

cuda-python (from cuda import cudart/cuda) wraps the CUDA driver directly. There is no drop-in ROCm binding; the driver calls must be rewritten against HIP.

RecommendRewrite driver-API calls against the HIP runtime.

FlashAttention dependency

from flash_attn.cute.flash_fwd_sm90 import FlashAttentionForwardSm90

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import (

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.compute_block_sparsity import compute_block_sparsity

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

compute_capability = torch.cuda.get_device_capability()

template:91

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

min_len, max_len + 1, (self.config.batch_size,), dtype=torch.int32, device="cuda"

template:142

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.zeros(1, dtype=torch.int32, device="cuda"),

template:146

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:156

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

current_stream = cuda.CUstream(torch.cuda.current_stream().cuda_stream)

template:338

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:534

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

start = torch.cuda.Event(enable_timing=True)

template:539

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

end = torch.cuda.Event(enable_timing=True)

template:540

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:545

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

benchmarks/bench_sm90.py· 15

FlashAttention dependency

from flash_attn.cute.interface import _flash_attn_fwd, _flash_attn_bwd

template:40

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

q = torch.randn(batch, seqlen, nheads, hdim, dtype=torch.bfloat16, device="cuda")

template:113

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch, seqlen, nheads, hdim, dtype=torch.bfloat16, device="cuda")

template:114

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch, seqlen, nheads, hdim_v, dtype=torch.bfloat16, device="cuda")

template:115

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:141

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

start = torch.cuda.Event(enable_timing=True)

template:142

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

end = torch.cuda.Event(enable_timing=True)

template:143

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:148

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

q = torch.randn(batch, seqlen, nheads, hdim, device="cuda", dtype=torch.bfloat16)

template:158

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch, seqlen, nheads, hdim, device="cuda", dtype=torch.bfloat16)

template:159

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch, seqlen, nheads, hdim_v, device="cuda", dtype=torch.bfloat16)

template:160

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:180

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

start = torch.cuda.Event(enable_timing=True)

template:181

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

end = torch.cuda.Event(enable_timing=True)

template:182

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:187

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

benchmarks/benchmark_alibi.py· 6

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.benchmark import benchmark_all, benchmark_forward, benchmark_backward

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.benchmark import benchmark_fwd_bwd, benchmark_combined

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn import flash_attn_qkvpacked_func, flash_attn_func

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

xFormers dependency

import xformers.ops as xops

template:18

xFormers bundles CUDA kernels. Many ops fall back to PyTorch SDPA on ROCm; some need a ROCm build.

RecommendPrefer torch SDPA; use a ROCm xFormers build where required.

Device string "cuda"

device = 'cuda'

template:116

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

benchmarks/benchmark_attn.py· 9

FlashAttention dependency

from flash_attn.cute.bench_utils import (

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.flash_attn_interface import flash_attn_func, flash_attn_varlen_func

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import flash_attn_func as flash_attn_func_python

template:27

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import flash_attn_varlen_func as flash_attn_varlen_func_python

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

if torch.cuda.get_device_capability()[0] != 9:

template:39

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Triton dependency

from triton.testing import do_bench

template:42

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

torch.cuda API usage

device_name = torch.cuda.get_device_name(device_index)

template:214

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

device_name = torch.cuda.get_device_name(device_index)

template:249

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = 'cuda'

template:397

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

benchmarks/benchmark_causal.py· 4

FlashAttention dependency

from flash_attn.utils.benchmark import benchmark_forward, benchmark_backward, benchmark_combined, benchmark_all, benchmark_fwd_bwd, pytorch_profiler

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.flash_attn_interface import flash_attn_varlen_qkvpacked_func

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn import flash_attn_qkvpacked_func, flash_attn_kvpacked_func

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = 'cuda'

template:64

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

benchmarks/benchmark_flash_attention.py· 6

FlashAttention dependency

from flash_attn.utils.benchmark import benchmark_all, benchmark_forward, benchmark_backward

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.benchmark import benchmark_fwd_bwd, benchmark_combined

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn import flash_attn_qkvpacked_func

template:14

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Triton dependency

from triton.ops.flash_attention import attention as attention_triton

template:17

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

xFormers dependency

import xformers.ops as xops

template:22

xFormers bundles CUDA kernels. Many ops fall back to PyTorch SDPA on ROCm; some need a ROCm build.

RecommendPrefer torch SDPA; use a ROCm xFormers build where required.

Device string "cuda"

device = 'cuda'

template:71

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

benchmarks/benchmark_gemm.py· 2

Triton dependency

from triton.testing import do_bench

template:5

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Device string "cuda"

device = 'cuda'

template:30

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

benchmarks/benchmark_mla_paged_kv.py· 3

Triton dependency

from triton.testing import do_bench

template:12

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

FlashAttention dependency

from flash_attn.cute.interface import flash_attn_varlen_func

template:14

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:17

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

benchmarks/clc_bench.py· 23

FlashAttention dependency

from flash_attn.cute import utils as cute_utils

template:376

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import flash_attn_func, flash_attn_varlen_func

template:377

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.compute_block_sparsity import compute_block_sparsity

template:389

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

cu_seqlens = torch_mod.zeros(len(lengths) + 1, device="cuda", dtype=torch_mod.int32)

template:396

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

cu_seqlens[1:] = torch_mod.tensor(lengths, device="cuda", dtype=torch_mod.int32).cumsum(0)

template:397

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = factory(case.batch, case.seqlen_q, case.q_heads, case.d, device="cuda", dtype=dtype)

template:402

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = factory(case.batch, case.seqlen_k, case.kv_heads, case.d, device="cuda", dtype=dtype)

template:403

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = factory(case.batch, case.seqlen_k, case.kv_heads, case.dv, device="cuda", dtype=dtype)

template:404

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = factory(total_q, case.q_heads, case.d, device="cuda", dtype=dtype)

template:413

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = factory(total_k, case.kv_heads, case.d, device="cuda", dtype=dtype)

template:414

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = factory(total_k, case.kv_heads, case.dv, device="cuda", dtype=dtype)

template:415

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

mask_block_cnt=torch_mod.zeros(count_shape, device="cuda", dtype=torch_mod.int32),

template:433

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

mask_block_idx=torch_mod.zeros(index_shape, device="cuda", dtype=torch_mod.int32),

template:434

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

full_block_cnt=torch_mod.zeros(count_shape, device="cuda", dtype=torch_mod.int32),

template:435

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

full_block_idx=torch_mod.zeros(index_shape, device="cuda", dtype=torch_mod.int32),

template:436

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = tensor_factory(case.batch, case.seqlen_q, case.q_heads, case.d, device="cuda", dtype=dtype)

template:445

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = tensor_factory(case.batch, case.seqlen_k, case.kv_heads, case.d, device="cuda", dtype=dtype)

template:446

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = tensor_factory(case.batch, case.seqlen_k, case.kv_heads, case.dv, device="cuda", dtype=dtype)

template:447

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:466

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if not fake_tensor and not torch.cuda.is_available():

template:640

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:648

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

if not torch.cuda.is_available():

template:707

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:719

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

csrc/flash_attn/src/flash_bwd_hdim128_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim128_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim128_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim128_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim192_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim192_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim192_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim192_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim256_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim256_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim256_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim256_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim32_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim32_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim32_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim32_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim64_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim64_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim64_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim64_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim96_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim96_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim96_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_bwd_hdim96_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim128_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim128_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim128_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim128_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim192_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim192_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim192_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim192_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim256_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim256_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim256_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim256_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim32_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim32_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim32_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim32_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim64_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim64_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim64_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim64_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim96_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim96_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim96_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_hdim96_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim128_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim128_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim128_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim128_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim192_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim192_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim192_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim192_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim256_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim256_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim256_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim256_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim32_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim32_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim32_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim32_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim64_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim64_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim64_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim64_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim96_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim96_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim96_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_align_hdim96_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim128_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim128_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim128_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim128_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim192_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim192_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim192_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim192_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim256_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim256_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim256_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim256_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim32_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim32_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim32_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim32_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim64_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim64_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim64_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim64_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim96_bf16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim96_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim96_fp16_causal_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/flash_fwd_split_hdim96_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/philox_unpack.cuh· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/flash_attn/src/philox.cuh· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/fused_dense_lib/fused_dense_cuda.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_1024.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_1280.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_1536.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_2048.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_256.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_2560.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_3072.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_4096.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_512.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_5120.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_6144.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_7168.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_768.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_8192.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_bwd_kernels.cuh· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_1024.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_1280.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_1536.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_2048.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_256.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_2560.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_3072.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_4096.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_512.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_5120.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_6144.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_7168.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_768.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_8192.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_fwd_kernels.cuh· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_1024.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_1280.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_1536.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_2048.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_256.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_2560.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_3072.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_4096.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_512.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_5120.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_6144.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_7168.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_768.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_bwd_8192.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_1024.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_1280.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_1536.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_2048.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_256.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_2560.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_3072.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_4096.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_512.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_5120.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_6144.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_7168.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_768.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_fwd_8192.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_residual_bwd_kernels.cuh· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_parallel_residual_fwd_kernels.cuh· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

csrc/layer_norm/ln_utils.cuh· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

flash_attn/__init__.py· 1

FlashAttention dependency

from flash_attn.flash_attn_interface import (

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/benchmark_flash_attention_fp8.py· 8

FlashAttention dependency

from flash_attn.cute.benchmark import benchmark_forward

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import _flash_attn_fwd as flash_attn_cute_fwd

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

default_scale_gpu = torch.ones(1, 1, 1, 1, dtype=torch.float32, device="cuda")

template:167

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

workspace = torch.empty(graph.get_workspace_size(), device="cuda", dtype=torch.uint8)

template:214

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if not torch.cuda.is_available():

template:253

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

major, minor = torch.cuda.get_device_capability()

template:255

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:262

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:276

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/cute/blackwell_helpers.py· 1

FlashAttention dependency

import flash_attn.cute.mma_sm100_desc as sm100_desc

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/block_info.py· 1

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK, SeqlenInfoQKNewK

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/block_sparse_utils.py· 4

FlashAttention dependency

from flash_attn.cute.block_sparsity import BlockSparseTensors

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.named_barrier import NamedBarrierBwd

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/block_sparsity.py· 1

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import get_broadcast_dims, to_cute_tensor

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/cache_utils.py· 1

FlashAttention dependency

from flash_attn.cute.fa_logging import fa_log

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/compute_block_sparsity.py· 8

FlashAttention dependency

from flash_attn.cute.block_sparsity import (

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparse_utils import get_curr_blocksparse_tensors

template:14

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.testing import is_fake_mode

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cute_dsl_utils import (

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import (

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.mask import call_mask_mod

template:27

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:29

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_bwd_sm120.py· 1

FlashAttention dependency

from flash_attn.cute.flash_bwd import FlashAttentionBackwardSm80

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/flash_fwd_sm120.py· 1

FlashAttention dependency

from flash_attn.cute.flash_fwd import FlashAttentionForwardSm80

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/mask.py· 4

FlashAttention dependency

import flash_attn.cute.utils as utils

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_info import BlockInfo

template:14

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/pack_gqa.py· 1

FlashAttention dependency

import flash_attn.cute.utils as utils

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/paged_kv.py· 1

FlashAttention dependency

from flash_attn.cute import utils

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/pyproject.toml· 5

FlashAttention dependency

name = "flash-attn-4"

template:6

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Pinned nvidia-* / +cuXXX wheel

"nvidia-cutlass-dsl==4.6.0.dev0",

template:25

A CUDA-pinned wheel (nvidia-* runtime package or a +cuXXX local version) forces the CUDA toolchain. Repin to the ROCm index.

RecommendRemove nvidia-* pins and install from the ROCm PyTorch index (--index-url .../whl/rocm6.1).

Pinned nvidia-* / +cuXXX wheel

cu13 = ["nvidia-cutlass-dsl[cu13]==4.6.0.dev0"]

template:35

A CUDA-pinned wheel (nvidia-* runtime package or a +cuXXX local version) forces the CUDA toolchain. Repin to the ROCm index.

RecommendRemove nvidia-* pins and install from the ROCm PyTorch index (--index-url .../whl/rocm6.1).

FlashAttention dependency

packages = ["flash_attn.cute"]

template:47

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

package-dir = {"flash_attn.cute" = "."}

template:48

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/softmax.py· 3

FlashAttention dependency

import flash_attn.cute.utils as utils

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.seqlen_info import SeqlenInfoQK

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import AuxData

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/tile_scheduler.py· 2

FlashAttention dependency

import flash_attn.cute.utils as utils

template:29

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.fast_math import clz

template:30

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/cute/topk_gather_kv.py· 2

FlashAttention dependency

from flash_attn.cute import utils

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import warp_reduce

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/flash_attn_interface.py· 11

FlashAttention dependency

import flash_attn_2_cuda

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

import flash_attn_2_cuda as flash_attn_gpu

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

major, minor = torch.cuda.get_device_capability(device)

template:34

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

@_torch_custom_op_wrapper("flash_attn::_flash_attn_forward", mutates_args=(), device_types="cuda")

template:84

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if torch.cuda.is_available() and torch.version.hip:

template:138

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

@_torch_custom_op_wrapper("flash_attn::_flash_attn_varlen_forward", mutates_args=(), device_types="cuda")

template:153

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if torch.cuda.is_available() and torch.version.hip:

template:238

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

@_torch_custom_op_wrapper("flash_attn::_flash_attn_backward", mutates_args=("dq", "dk", "dv"), device_types="cuda")

template:252

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if torch.cuda.is_available() and torch.version.hip:

template:333

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

@_torch_custom_op_wrapper("flash_attn::_flash_attn_varlen_backward", mutates_args=("dq", "dk", "dv"), device_types="cuda")

template:347

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if torch.cuda.is_available() and torch.version.hip:

template:447

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/flash_attn_triton_og.py· 2

Triton dependency

import triton

template:14

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Triton dependency

import triton.language as tl

template:15

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

flash_attn/flash_attn_triton.py· 2

Triton dependency

import triton

template:45

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Triton dependency

import triton.language as tl

template:46

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

flash_attn/flash_blocksparse_attention.py· 2

FlashAttention dependency

from flash_attn.bert_padding import index_first_axis, pad_input, unpad_input

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.flash_blocksparse_attn_interface import (

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/layers/patch_embed.py· 1

FlashAttention dependency

from flash_attn.ops.fused_dense import FusedDense

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/layers/rotary.py· 1

FlashAttention dependency

from flash_attn.ops.triton.rotary import apply_rotary

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/losses/cross_entropy.py· 1

FlashAttention dependency

from flash_attn.ops.triton.cross_entropy import cross_entropy_loss

template:6

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/models/bert.py· 9

FlashAttention dependency

from flash_attn.bert_padding import (

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.block import Block

template:31

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.embedding import BertEmbeddings

template:32

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.mha import MHA

template:33

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.mlp import FusedMLP, Mlp

template:34

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:35

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.fused_dense import FusedDense

template:38

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.triton.layer_norm import layer_norm_fn

template:43

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.losses.cross_entropy import CrossEntropyLoss

template:49

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/models/gpt.py· 17

FlashAttention dependency

from flash_attn.models.bigcode import remap_state_dict_hf_bigcode

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.falcon import remap_state_dict_hf_falcon

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.gpt_neox import remap_state_dict_hf_gpt_neox

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.gptj import remap_state_dict_hf_gptj

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.llama import remap_state_dict_hf_llama

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.opt import remap_state_dict_hf_opt

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.block import Block, ParallelBlock

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.embedding import GPT2Embeddings, ParallelGPT2Embeddings

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.mha import MHA, ParallelMHA

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.mlp import (

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.activations import sqrelu_fwd

template:34

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.distributed import (

template:35

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.generation import GenerationMixin

template:41

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:42

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.fused_dense import ColumnParallelLinear

template:45

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.triton.mlp import FusedDenseSqreluDense

template:50

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.triton.layer_norm import layer_norm_fn, RMSNorm

template:55

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/models/vit.py· 5

FlashAttention dependency

from flash_attn.layers.patch_embed import PatchEmbed

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.block import Block

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.mha import MHA

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.mlp import FusedMLP, Mlp

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.triton.layer_norm import layer_norm_fn

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/modules/block.py· 3

FlashAttention dependency

from flash_attn.modules.mha import MHA

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.mlp import Mlp

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.triton.layer_norm import layer_norm_fn, RMSNorm

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/modules/embedding.py· 1

FlashAttention dependency

from flash_attn.utils.distributed import all_reduce, reduce_scatter

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/modules/mha.py· 4

FlashAttention dependency

from flash_attn.utils.distributed import get_dim_for_local_rank

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn import (

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.fused_dense import ColumnParallelLinear, RowParallelLinear

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.layers.rotary import RotaryEmbedding

template:31

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/modules/mlp.py· 3

FlashAttention dependency

from flash_attn.ops.activations import swiglu

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.fused_dense import ColumnParallelLinear, RowParallelLinear

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.fused_dense import FusedMLP, ParallelFusedMLP

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/ops/fused_dense.py· 5

FlashAttention dependency

from flash_attn.utils.torch import custom_fwd, custom_bwd

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.activations import gelu_bwd, relu_bwd, sqrelu_bwd, sqrelu_fwd

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.distributed import (

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

if torch.cuda.get_device_capability("cuda") == (9, 0):

template:584

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if torch.cuda.get_device_capability("cuda") == (9, 0):

template:584

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/ops/rms_norm.py· 1

FlashAttention dependency

from flash_attn.ops.layer_norm import (

template:7

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/ops/triton/cross_entropy.py· 4

Triton dependency

import triton

template:8

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Triton dependency

import triton.language as tl

template:9

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

torch.cuda API usage

with torch.cuda.device(logits.device.index):

template:196

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

with torch.cuda.device(logits.device.index):

template:269

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/ops/triton/k_activations.py· 2

Triton dependency

import triton

template:11

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Triton dependency

import triton.language as tl

template:12

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

flash_attn/ops/triton/layer_norm.py· 10

Triton dependency

import triton

template:16

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Triton dependency

import triton.language as tl

template:17

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

FlashAttention dependency

from flash_attn.utils.torch import custom_fwd, custom_bwd

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.library import triton_op

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

warp_size = getattr(torch.cuda.get_device_properties(torch.cuda.current_device()), "warp_size", 32)

template:37

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

with torch.cuda.device(x.device.index):

template:426

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

sm_count = torch.cuda.get_device_properties(x.device).multi_processor_count * 8

template:779

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

with torch.cuda.device(x.device.index):

template:790

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

out_dtype=None if not torch.is_autocast_enabled() else torch.get_autocast_dtype("cuda"),

template:1165

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

dtype = torch.get_autocast_dtype("cuda") if torch.is_autocast_enabled() else y.dtype

template:1170

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

flash_attn/ops/triton/linear.py· 4

Triton dependency

import triton

template:6

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Triton dependency

import triton.language as tl

template:7

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Triton dependency

from triton.ops.matmul_perf_model import early_config_prune, estimate_matmul_time

template:8

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

FlashAttention dependency

from flash_attn.ops.triton.k_activations import (

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/ops/triton/mlp.py· 3

FlashAttention dependency

from flash_attn.utils.torch import custom_fwd, custom_bwd

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.activations import sqrelu_bwd, sqrelu_fwd

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.triton.linear import triton_dgrad_act, triton_linear_act

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

flash_attn/ops/triton/rotary.py· 3

Triton dependency

import triton

template:8

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Triton dependency

import triton.language as tl

template:9

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

torch.cuda API usage

with torch.cuda.device(x.device.index):

template:158

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/utils/testing.py· 1

FlashAttention dependency

from flash_attn.bert_padding import pad_input, unpad_input

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

hopper/benchmark_attn.py· 7

FlashAttention dependency

from flash_attn.utils.benchmark import benchmark_forward, benchmark_backward, benchmark_combined, benchmark_all, benchmark_fwd_bwd, pytorch_profiler

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.flash_attn_interface import flash_attn_func, flash_attn_varlen_func

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Triton dependency

from triton.testing import do_bench

template:30

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Device string "cuda"

row_idx = torch.arange(seqlen_q, device='cuda')

template:69

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

workspace = torch.empty(graph.get_workspace_size(), device="cuda", dtype=torch.uint8)

template:137

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

workspace = torch.empty(graph.get_workspace_size(), device="cuda", dtype=torch.uint8)

template:206

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = 'cuda'

template:222

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

hopper/benchmark_flash_attention_fp8.py· 11

FlashAttention dependency

from flash_attn.utils.benchmark import benchmark_all, benchmark_forward, benchmark_backward

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.benchmark import benchmark_fwd_bwd, benchmark_combined

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn import flash_attn_qkvpacked_func

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

xFormers dependency

import xformers.ops as xops

template:24

xFormers bundles CUDA kernels. Many ops fall back to PyTorch SDPA on ROCm; some need a ROCm build.

RecommendPrefer torch SDPA; use a ROCm xFormers build where required.

Device string "cuda"

default_scale_gpu = torch.ones(1, 1, 1, 1, dtype=torch.float32, device="cuda")

template:113

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

workspace = torch.empty(graph.get_workspace_size(), device="cuda", dtype=torch.uint8)

template:164

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = 'cuda'

template:218

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:248

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

descale_q = torch.tensor([1.0], dtype=torch.float32, device='cuda')

template:282

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

descale_k = torch.tensor([1.0], dtype=torch.float32, device='cuda')

template:283

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

descale_v = torch.tensor([1.0], dtype=torch.float32, device='cuda')

template:284

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

hopper/benchmark_mla_decode.py· 7

Triton dependency

from triton.testing import do_bench, do_bench_cudagraph

template:13

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

FlashAttention dependency

from flash_attn.utils.benchmark import pytorch_profiler

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:30

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize() # Gotta wait, otherwise e.g. k_cache might not be ready

template:94

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

with torch.cuda.stream(torch.cuda.Stream()):

template:95

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize() # Gotta wait, otherwise e.g. k_cache might not be ready

template:108

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

with torch.cuda.stream(torch.cuda.Stream()):

template:109

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

hopper/benchmark_split_kv.py· 9

FlashAttention dependency

import flash_attn

template:2

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

torch.cuda.synchronize()

template:16

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

num_sms = torch.cuda.get_device_properties(

template:36

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.current_device()

template:37

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=dtype

template:91

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=dtype

template:94

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn((num_requests, query_seqlen, nheads_q, headdim), device="cuda", dtype=dtype)

template:108

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

cache_idxs = torch.randperm(num_caches, dtype=torch.int32, device="cuda")[:num_requests]

template:109

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

[context_seqlen] * num_requests, dtype=torch.int32, device="cuda"

template:111

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

hopper/flash_fwd_combine.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/flash_prepare_scheduler.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_bf16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim128_fp16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_bf16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim192_fp16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_bf16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim256_fp16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_bf16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim64_fp16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_bf16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_bwd_hdim96_fp16_softcapall_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_sm100.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_bf16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_e4m3_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim128_fp16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_e4m3_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_128_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_bf16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_e4m3_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim192_fp16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_bf16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_e4m3_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim256_fp16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_256_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_512_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_bf16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_e4m3_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim64_fp16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_bf16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_e4m3_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_paged_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_split_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_split_softcap_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdim96_fp16_split_softcapall_sm80.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_e4m3_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimall_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_bf16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_e4m3_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_paged_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_paged_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_paged_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_paged_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_softcap_packgqa_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_split_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/instantiations/flash_fwd_hdimdiff_fp16_split_softcap_sm90.cu· 1

Custom CUDA kernels (.cu/.cuh)

CUDA kernel source

templatefile

Hand-written .cu/.cuh kernels compile only with nvcc. AMD's HIPify tools translate the vast majority of CUDA C++ to HIP automatically.

RecommendRun hipify-perl over the .cu/.cuh sources and build with hipcc.

hopper/test_attn_kvcache.py· 37

FlashAttention dependency

import flash_attn

template:4

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:156

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=torch.bfloat16

template:162

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=torch.bfloat16

template:165

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn((num_requests, query_seqlen, nheads_q, headdim), device="cuda", dtype=torch.bfloat16)

template:167

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

cache_seqlens = torch.tensor([context_seqlen] * num_requests, dtype=torch.int32, device="cuda")

template:169

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:170

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:192

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:218

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=torch.bfloat16

template:224

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=torch.bfloat16

template:227

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn((num_requests, query_seqlen, nheads_q, headdim), device="cuda", dtype=torch.bfloat16)

template:229

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

cache_seqlens = torch.tensor([context_seqlen] * num_requests, dtype=torch.int32, device="cuda")

template:234

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:235

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

descale_q = torch.tensor([1.0], dtype=torch.float32, device='cuda')

template:244

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

descale_k = torch.tensor([1.0], dtype=torch.float32, device='cuda')

template:245

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

descale_v = torch.tensor([1.0], dtype=torch.float32, device='cuda')

template:246

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:261

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:293

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=torch.bfloat16

template:306

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=torch.bfloat16

template:309

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn((num_requests, query_seqlen, nheads_q, headdim), device="cuda", dtype=torch.bfloat16)

template:311

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

cache_idxs = torch.randperm(num_caches, dtype=torch.int32, device="cuda")[:num_requests]

template:316

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

cache_seqlens = torch.randint(1, context_seqlen-1, (num_requests,), dtype=torch.int32).to(device) if cache_seqlen_rand else torch.tensor([context_seqlen] * num_requests, dtype=torch.int32, device="cud

template:317

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:318

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:347

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:400

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=torch.bfloat16

template:413

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=torch.bfloat16

template:416

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn((num_requests, query_seqlen, nheads_q, headdim), device="cuda", dtype=torch.bfloat16)

template:418

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

cache_idxs = torch.randperm(num_caches, dtype=torch.int32, device="cuda")[:num_requests]

template:423

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

template:424

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:425

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

descale_q = torch.tensor([1.0], dtype=torch.float32, device='cuda')

template:428

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

descale_k = torch.tensor([1.0], dtype=torch.float32, device='cuda')

template:429

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

descale_v = torch.tensor([1.0], dtype=torch.float32, device='cuda')

template:430

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:461

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

hopper/test_flash_attn_bwd_determinism.py· 5

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

DISABLE_FP8 = os.getenv("FLASH_ATTENTION_DISABLE_FP8", "FALSE") == "TRUE" or torch.cuda.get_device_capability("cuda")[0] < 9

template:37

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

DISABLE_FP8 = os.getenv("FLASH_ATTENTION_DISABLE_FP8", "FALSE") == "TRUE" or torch.cuda.get_device_capability("cuda")[0] < 9

template:37

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:119

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:398

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

hopper/test_flash_attn_triton_amd.py· 9

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

DISABLE_FP8 = os.getenv("FLASH_ATTENTION_DISABLE_FP8", "FALSE") == "TRUE" or torch.cuda.get_device_capability("cuda")[0] < 9

template:35

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

DISABLE_FP8 = os.getenv("FLASH_ATTENTION_DISABLE_FP8", "FALSE") == "TRUE" or torch.cuda.get_device_capability("cuda")[0] < 9

template:35

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:110

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:337

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:654

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:999

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1043

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1102

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

hopper/test_flash_attn.py· 9

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

DISABLE_FP8 = os.getenv("FLASH_ATTENTION_DISABLE_FP8", "FALSE") == "TRUE" or torch.cuda.get_device_capability("cuda")[0] < 9

template:40

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

DISABLE_FP8 = os.getenv("FLASH_ATTENTION_DISABLE_FP8", "FALSE") == "TRUE" or torch.cuda.get_device_capability("cuda")[0] < 9

template:40

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:174

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:409

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:741

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1091

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1134

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1193

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

hopper/test_kvcache.py· 11

FlashAttention dependency

import flash_attn as fa2

template:4

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

torch.cuda.synchronize()

template:39

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:43

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=dtype

template:83

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(num_caches, cache_seqlen, nheads_kv, headdim), device="cuda", dtype=dtype

template:86

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

(1, ntokens, nheads_q, headdim), device="cuda", dtype=dtype

template:90

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

[context_seqlens[0]], dtype=torch.int32, device="cuda"

template:93

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

cache_idx_large = torch.tensor([1], dtype=torch.int32, device="cuda")

template:95

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:99

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

context_seqlens[1:] + [0] * num_padding_queries, dtype=torch.int32, device="cuda"

template:103

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

cache_idxs_small = torch.randperm(num_caches, dtype=torch.int32, device="cuda")[

template:105

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/benchmark_block_sparsity.py· 8

FlashAttention dependency

from flash_attn.cute.compute_block_sparsity import BlockSparsityKernel

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import BlockSparseTensors

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Triton dependency

from triton.testing import do_bench

template:25

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

Device string "cuda"

device = "cuda"

template:64

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:99

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize(device)

template:185

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:311

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:386

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/cute/mask_mod_definitions.py· 6

FlashAttention dependency

from flash_attn.cute import utils

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import fast_sampling

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

def make_packed_doc_ids(seqlens_q, seqlens_k, device="cuda"):

template:422

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

def make_global_thresholds(seqlens_k, device="cuda"):

template:466

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

def make_global_windows(seqlens_q, device="cuda"):

template:483

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:694

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_block_sparsity.py· 9

FlashAttention dependency

from flash_attn.cute.compute_block_sparsity import compute_block_sparsity

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device="cuda",

template:36

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:239

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:321

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:388

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:451

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:520

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:621

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:690

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_cache_utils.py· 2

FlashAttention dependency

import flash_attn.cute.cache_utils as cache_utils

template:4

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import fa_logging

template:5

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

tests/cute/test_clc_fuzz.py· 28

FlashAttention dependency

from flash_attn.cute import utils as cute_utils

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.flash_fwd_sm100 import FlashAttentionForwardSm100

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import flash_attn_func, flash_attn_varlen_func

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.testing import attention_ref

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import (

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

if torch.cuda.is_available():

template:28

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

COMPUTE_CAPABILITY = torch.cuda.get_device_capability()[0]

template:29

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

SM_COUNT = torch.cuda.get_device_properties("cuda").multi_processor_count

template:30

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

SM_COUNT = torch.cuda.get_device_properties("cuda").multi_processor_count

template:30

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:65

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

return torch.randn(b, s, h, d, device="cuda", dtype=torch.bfloat16)

template:96

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

FlashAttention dependency

from flash_attn.cute.tile_scheduler import SingleTileScheduler

template:254

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

cu_seqlens = torch.cat([torch.zeros(1, dtype=torch.int32), lens.cumsum(0)]).to(device="cuda", dtype=torch.int32)

template:270

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn(total, heads, d, device="cuda", dtype=torch.bfloat16)

template:272

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(total, heads, d, device="cuda", dtype=torch.bfloat16)

template:273

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(total, heads, d, device="cuda", dtype=torch.bfloat16)

template:274

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:282

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

cu_seqlens = torch.cat([torch.zeros(1, dtype=torch.int32), torch.tensor(seqlens, dtype=torch.int32).cumsum(0)]).to(device="cuda", dtype=torch.int32)

template:307

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn(total, heads, d, device="cuda", dtype=torch.bfloat16)

template:310

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(total, kv_heads, d, device="cuda", dtype=torch.bfloat16)

template:311

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(total, kv_heads, d, device="cuda", dtype=torch.bfloat16)

template:312

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:324

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

seqused = torch.tensor(seqlens, device="cuda", dtype=torch.int32)

template:352

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn(batch, max_s, heads, d, device="cuda", dtype=torch.bfloat16)

template:353

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch, max_s, kv_heads, d, device="cuda", dtype=torch.bfloat16)

template:354

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch, max_s, kv_heads, d, device="cuda", dtype=torch.bfloat16)

template:355

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q_mask = torch.arange(max_s, device="cuda")[None, :] < seqused[:, None]

template:356

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:371

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/cute/test_flash_attn_combine.py· 5

FlashAttention dependency

from flash_attn.cute.testing import (

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import (

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:59

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:116

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:240

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_flash_attn_fast.py· 9

FlashAttention dependency

from flash_attn.cute.testing import (

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import (

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

IS_SM90 = torch.cuda.get_device_capability()[0] == 9

template:27

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

IS_SM120 = torch.cuda.get_device_capability()[0] == 12

template:28

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:53

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:56

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:119

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:192

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:310

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_flash_attn_race_condition.py· 10

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.testing import (

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import (

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

IS_SM90 = torch.cuda.get_device_capability()[0] == 9

template:33

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

is_sm90 = torch.cuda.get_device_capability()[0] == 9

template:79

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:82

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:85

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:86

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

is_sm90 = torch.cuda.get_device_capability()[0] == 9

template:411

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:420

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_flash_attn_varlen.py· 2

FlashAttention dependency

from flash_attn.cute import flash_attn_varlen_func

template:6

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:158

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_flash_attn.py· 39

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.testing import (

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import (

template:30

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

torch.cuda.empty_cache()

template:49

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

IS_SM90 = torch.cuda.get_device_capability()[0] == 9

template:83

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

IS_SM100 = torch.cuda.get_device_capability()[0] == 10

template:84

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

IS_SM120 = torch.cuda.get_device_capability()[0] == 12

template:85

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

q = torch.randn(1, 16, 4, 64, device="cuda", dtype=torch.bfloat16)

template:92

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(1, 16, 1, 64, device="cuda", dtype=torch.bfloat16)

template:93

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(1, 16, 1, 64, device="cuda", dtype=torch.bfloat16)

template:94

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:192

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:197

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:198

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:485

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:489

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:490

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

q = torch.randn(batch_size, nheads, seqlen, d, device="cuda", dtype=dtype).transpose(1, 2)

template:526

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch_size, nheads, seqlen, d, device="cuda", dtype=dtype).transpose(1, 2)

template:527

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch_size, nheads, seqlen, d, device="cuda", dtype=dtype).transpose(1, 2)

template:528

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:660

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1155

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

FlashAttention dependency

from flash_attn.cute.interface import _flash_attn_fwd, _flash_attn_bwd

template:1638

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:1640

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1678

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1763

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1847

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1900

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1972

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2034

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2093

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2162

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:2167

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:2168

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:2396

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2758

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2859

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2916

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2954

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:3004

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_mask_mod_varlen.py· 14

FlashAttention dependency

from flash_attn.cute.interface import _flash_attn_fwd

template:21

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import utils

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.compute_block_sparsity import compute_block_sparsity

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

COMPUTE_CAPABILITY = torch.cuda.get_device_capability()[0]

template:43

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.empty_cache()

template:50

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.empty_cache()

template:53

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:104

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

num_heads, batch_size, max_doc_len, device="cuda"

template:292

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

).to(dtype=torch.int32, device="cuda")

template:293

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

doc_ids_q, doc_ids_k = make_packed_doc_ids(seqlens_q, seqlens_k, device="cuda")

template:519

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

thresholds = make_global_thresholds(seqlens_k, device="cuda")

template:526

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

windows = make_global_windows(seqlens_q, device="cuda")

template:531

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:916

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1088

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_mask_mod.py· 71

FlashAttention dependency

from flash_attn.cute.interface import _flash_attn_fwd, _flash_attn_bwd, flash_attn_func

template:25

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import (

template:26

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.cache_utils import get_jit_cache

template:33

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.compute_block_sparsity import compute_block_sparsity

template:34

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute import utils

template:35

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

COMPUTE_CAPABILITY = torch.cuda.get_device_capability()[0]

template:43

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.empty_cache()

template:62

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.empty_cache()

template:67

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:73

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

doc_ids = random_doc_id_tensor(nheads, batch_size, doc_len, device="cuda").to(

template:315

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

dtype=torch.int32, device="cuda"

template:316

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

bias = torch.full((seqlen_k,), bias_threshold, dtype=torch.int32, device="cuda")

template:327

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:356

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:383

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda", BLOCK_SIZE=(tile_m, tile_n),

template:532

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

doc_ids = torch.zeros(batch_size, nheads, max(seqlen_q, seqlen_k), dtype=torch.int32, device="cuda")

template:660

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn(batch_size, seqlen_q, nheads, headdim, device="cuda", dtype=dtype)

template:672

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch_size, seqlen_k, nheads, headdim, device="cuda", dtype=dtype)

template:673

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch_size, seqlen_k, nheads, headdim, device="cuda", dtype=dtype)

template:674

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

out = torch.empty(batch_size, seqlen_q, nheads, headdim, device="cuda", dtype=dtype)

template:675

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

lse = torch.empty(batch_size, nheads, seqlen_q, device="cuda", dtype=torch.float32)

template:676

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda", BLOCK_SIZE=(sparse_tile_m, tile_n),

template:681

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda", BLOCK_SIZE=(tile_m, tile_n),

template:736

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_size, seqlen_q, num_heads, head_dim, device="cuda", dtype=dtype

template:925

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_size, seqlen_k, num_kv_heads, head_dim, device="cuda", dtype=dtype

template:928

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_size, seqlen_k, num_kv_heads, head_dim, device="cuda", dtype=dtype

template:931

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if torch.cuda.get_device_capability()[0] != 10:

template:969

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:971

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

FlashAttention dependency

from flash_attn.cute import flash_fwd_sm100

template:1014

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.interface import _flash_attn_fwd

template:1015

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device="cuda",

template:1080

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:1177

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:1365

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

out = torch.empty(batch_size, seqlen_q, nheads, headdim, device="cuda", dtype=dtype)

template:1391

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

lse = torch.empty(batch_size, nheads, seqlen_q, device="cuda", dtype=torch.float32)

template:1392

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn(batch_size, seqlen_q, nheads, headdim, device="cuda", dtype=dtype)

template:1429

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch_size, seqlen_k, nheads, headdim, device="cuda", dtype=dtype)

template:1430

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch_size, seqlen_k, nheads, headdim, device="cuda", dtype=dtype)

template:1431

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:1442

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn(batch_size, seqlen_q, nheads, headdim, device="cuda", dtype=dtype)

template:1531

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch_size, seqlen_k, nheads, headdim, device="cuda", dtype=dtype)

template:1532

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch_size, seqlen_k, nheads, headdim, device="cuda", dtype=dtype)

template:1533

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:1544

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda", BLOCK_SIZE=(sparse_tile_m, tile_n),

template:1657

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda", BLOCK_SIZE=(tile_m, tile_n),

template:1701

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1742

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda", BLOCK_SIZE=(sparse_tile_m, tile_n),

template:1856

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn(batch_size, seqlen_q, nheads_q, headdim, device="cuda", dtype=dtype)

template:1865

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch_size, seqlen_k, nheads_kv, headdim, device="cuda", dtype=dtype)

template:1866

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch_size, seqlen_k, nheads_kv, headdim, device="cuda", dtype=dtype)

template:1867

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

out=torch.empty(batch_size, seqlen_q, nheads_q, headdim, device="cuda", dtype=dtype),

template:1871

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

lse=torch.empty(batch_size, nheads_q, seqlen_q, device="cuda", dtype=torch.float32),

template:1872

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:1884

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device="cuda", BLOCK_SIZE=(block_size, block_size),

template:2003

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:2047

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

doc_ids = random_doc_id_tensor(nheads, batch_size, max(seqlen_q, seqlen_k), device="cuda").to(

template:2121

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

out=torch.empty(batch_size, seqlen_q, nheads, headdim, device="cuda", dtype=dtype),

template:2155

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

lse=torch.empty(batch_size, nheads, seqlen_q, device="cuda", dtype=torch.float32),

template:2156

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn(batch_size, seqlen_q, nheads, headdim, device="cuda", dtype=dtype)

template:2237

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch_size, seqlen_k, nheads, headdim, device="cuda", dtype=dtype)

template:2238

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch_size, seqlen_k, nheads, headdim, device="cuda", dtype=dtype)

template:2239

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

out=torch.empty(batch_size, seqlen_q, nheads, headdim, device="cuda", dtype=dtype),

template:2254

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

lse=torch.empty(batch_size, nheads, seqlen_q, device="cuda", dtype=torch.float32),

template:2255

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

doc_ids = random_doc_id_tensor(H, B, max(seqlen_q, seqlen_k), device="cuda").to(torch.int32)

template:2389

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:2420

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:2493

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda", BLOCK_SIZE=(sparse_tile_m, tile_n),

template:2566

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn(1, 128, 4, 64, device="cuda", dtype=torch.bfloat16),

template:2645

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn(1, 128, 4, 64, device="cuda", dtype=torch.bfloat16),

template:2646

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn(1, 128, 4, 64, device="cuda", dtype=torch.bfloat16),

template:2647

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:2675

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_score_mod_varlen.py· 39

FlashAttention dependency

from flash_attn.cute.interface import _flash_attn_fwd

template:4

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

IS_SM90 = torch.cuda.get_device_capability()[0] == 9

template:67

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

IS_SM100 = torch.cuda.get_device_capability()[0] == 10

template:68

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

q = torch.randn(total_q, num_heads, head_dim, device="cuda", dtype=dtype)

template:289

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:292

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_size, seqlen_q, num_heads, head_dim, device="cuda", dtype=dtype

template:298

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(total_k, num_heads, head_dim, device="cuda", dtype=dtype)

template:304

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(total_k, num_heads, head_dim, device="cuda", dtype=dtype)

template:305

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:308

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_size, seqlen_k, num_heads, head_dim, device="cuda", dtype=dtype

template:314

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_size, seqlen_k, num_heads, head_dim, device="cuda", dtype=dtype

template:317

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:335

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

bias = torch.zeros(batch_size, device="cuda", dtype=dtype) * 0.1

template:460

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

head_bias = torch.randn(num_heads, device="cuda", dtype=dtype) * 0.2

template:465

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

pos_bias = torch.arange(seqlen_q, device="cuda", dtype=dtype) * 0.01

template:466

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

bias = torch.zeros(batch_size, device="cuda", dtype=dtype) * 0.1

template:555

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

head_bias = torch.randn(num_heads, device="cuda", dtype=dtype) * 0.2

template:559

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

pos_bias = torch.arange(seqlen_q, device="cuda", dtype=dtype) * 0.01

template:560

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:651

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device="cuda",

template:656

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn(total_q, num_heads, head_dim, device="cuda", dtype=dtype)

template:661

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_size, seqlen_q, num_heads, head_dim, device="cuda", dtype=dtype

template:665

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(total_k, num_heads, head_dim, device="cuda", dtype=dtype)

template:669

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(total_k, num_heads, head_dim, device="cuda", dtype=dtype)

template:670

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_size, seqlen_k, num_heads, head_dim, device="cuda", dtype=dtype

template:674

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_size, seqlen_k, num_heads, head_dim, device="cuda", dtype=dtype

template:677

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

bias = torch.randn(total_k, device="cuda", dtype=dtype) * 0.1

template:690

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

bias = torch.randn(total_q, device="cuda", dtype=dtype) * 0.1

template:694

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q_bias = torch.randn(total_q, device="cuda", dtype=dtype) * 0.1

template:698

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

kv_bias = torch.randn(total_k, device="cuda", dtype=dtype) * 0.1

template:699

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

bias = torch.randn(total_q, device="cuda", dtype=dtype) * 0.1

template:703

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

kv_bias = torch.randn(total_k, device="cuda", dtype=dtype) * 0.1

template:707

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_bias = torch.randn(batch_size, device="cuda", dtype=dtype) * 0.1

template:711

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

head_scale = torch.randn(num_heads, device="cuda", dtype=dtype) * 0.1 + 1.0

template:712

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q_pos_bias = torch.randn(total_q, device="cuda", dtype=dtype) * 0.1

template:713

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

kv_pos_bias = torch.randn(total_k, device="cuda", dtype=dtype) * 0.1

template:714

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn(max_rel_pos * 2 + 1, device="cuda", dtype=dtype) * 0.1

template:716

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:831

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:998

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_score_mod.py· 19

FlashAttention dependency

from flash_attn.cute.interface import (

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.block_sparsity import BlockSparseTensorsTorch

template:14

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

COMPUTE_CAPABILITY = torch.cuda.get_device_capability()[0]

template:16

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

COMPUTE_CAPABILITY = torch.cuda.get_device_capability()[0]

template:57

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

q = torch.randn(batch_size, num_heads, seqlen_q, dim, device="cuda", dtype=dtype)

template:160

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

k = torch.randn(batch_size, num_heads, seqlen_kv, dim, device="cuda", dtype=dtype)

template:161

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

v = torch.randn(batch_size, num_heads, seqlen_kv, dim, device="cuda", dtype=dtype)

template:162

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

buffer = torch.randn(batch_size, device="cuda", dtype=dtype) * 0.1

template:303

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

head_bias = torch.randn(num_q_heads, device="cuda", dtype=dtype) * 0.2

template:308

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

pos_scale = torch.arange(seqlen_q, device="cuda", dtype=dtype) * 0.01

template:309

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

buffer = torch.randn(batch_size, device="cuda", dtype=dtype) * 0.1

template:376

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

head_bias = torch.randn(num_q_heads, device="cuda", dtype=dtype) * 0.2

template:380

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

pos_scale = torch.arange(seqlen_q, device="cuda", dtype=dtype) * 0.01

template:381

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:464

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:613

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

buffer = torch.randn(batch_size, device="cuda", dtype=dtype) * 0.1

template:1104

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

head_bias = torch.randn(num_heads, device="cuda", dtype=dtype) * 0.2

template:1106

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

pos_scale = torch.arange(seqlen_q, device="cuda", dtype=dtype) * 0.01

template:1107

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

batch_bias = torch.randn(2, device="cuda", dtype=torch.bfloat16) * 0.1

template:1262

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/cute/test_utils.py· 2

FlashAttention dependency

from flash_attn.cute import utils as cute_utils

template:5

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.cute.utils import hash_callable

template:6

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

tests/layers/test_rotary.py· 3

FlashAttention dependency

from flash_attn.layers.rotary import RotaryEmbedding, apply_rotary_emb_func, apply_rotary_emb_qkv_

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:22

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:96

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/losses/test_cross_entropy_parallel.py· 4

NVIDIA Apex dependency

from apex.transformer import parallel_state, tensor_parallel

template:8

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

FlashAttention dependency

from flash_attn.losses.cross_entropy import CrossEntropyLoss

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:11

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:11

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/losses/test_cross_entropy.py· 4

FlashAttention dependency

from flash_attn.losses.cross_entropy import CrossEntropyLoss

template:6

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:8

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:8

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:40

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/models/test_baichuan.py· 18

FlashAttention dependency

from flash_attn.models.gpt import (

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.baichuan import (

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.distributed import all_gather_raw

template:22

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:23

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.generation import update_graph_cache

template:24

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:66

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:149

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

Device string "cuda"

device = "cuda"

template:235

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:261

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:269

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:291

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:302

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:309

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:320

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:350

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

torch.cuda API usage

torch.cuda.set_device(device)

template:383

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:430

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:439

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/models/test_bert.py· 18

FlashAttention dependency

from flash_attn.models.bert import (

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

.cuda() tensor/module move

model_hf.cuda().to(dtype=dtype)

template:48

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

.cuda() tensor/module move

model = model.cuda().to(dtype=dtype)

template:62

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

seqlens = torch.randint(max_seqlen // 2, max_seqlen + 1, (batch_size,), device="cuda")

template:74

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

attention_mask = torch.arange(max_seqlen, device="cuda")[None, :] < seqlens[:, None]

template:75

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

0, config.vocab_size, (batch_size, max_seqlen), dtype=torch.long, device="cuda"

template:77

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

.cuda() tensor/module move

model = model.cuda().to(dtype=dtype)

template:117

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

seqlens = torch.randint(max_seqlen // 2, max_seqlen + 1, (batch_size,), device="cuda")

template:129

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

attention_mask = torch.arange(max_seqlen, device="cuda")[None, :] < seqlens[:, None]

template:130

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

0, config.vocab_size, (batch_size, max_seqlen), dtype=torch.long, device="cuda"

template:132

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

.cuda() tensor/module move

model = model.cuda().to(dtype=dtype)

template:227

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

seqlens = torch.randint(max_seqlen // 2, max_seqlen + 1, (batch_size,), device="cuda")

template:239

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

attention_mask = torch.arange(max_seqlen, device="cuda")[None, :] < seqlens[:, None]

template:241

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

0, config.vocab_size, (batch_size, max_seqlen), dtype=torch.long, device="cuda"

template:245

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

0, config.vocab_size, (batch_size, max_seqlen), dtype=torch.long, device="cuda"

template:248

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

labels[(torch.rand(batch_size, max_seqlen, device="cuda") > 0.15)] = 0

template:252

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

next_sequence_label = torch.randint(0, 2, (batch_size,), device="cuda")

template:254

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/models/test_bigcode.py· 12

FlashAttention dependency

from flash_attn.models.bigcode import bigcode_config_to_gpt2_config, inv_remap_state_dict_hf_bigcode

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.gpt import GPTLMHeadModel, remap_state_dict_hf_bigcode

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.generation import update_graph_cache

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:34

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:94

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:119

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:124

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:138

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:149

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:156

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:167

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/models/test_btlm.py· 13

FlashAttention dependency

from flash_attn.models.gpt import GPTLMHeadModel

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.btlm import btlm_config_to_gpt2_config, remap_state_dict_hf_btlm

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.generation import update_graph_cache

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:36

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:102

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:126

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:134

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:154

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:165

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:172

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:183

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:208

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/models/test_falcon.py· 18

FlashAttention dependency

from flash_attn.models.falcon import falcon_config_to_gpt2_config, remap_state_dict_hf_falcon

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.gpt import GPTLMHeadModel, combine_state_dicts_tp, shard_state_dict_tp

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.distributed import all_gather_raw

template:14

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.generation import update_graph_cache

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:42

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:105

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

Device string "cuda"

device = "cuda"

template:192

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:218

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:223

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:239

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:250

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:257

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:268

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:299

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

torch.cuda API usage

torch.cuda.set_device(device)

template:332

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:378

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:387

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/models/test_gpt_generation_parallel.py· 5

FlashAttention dependency

from flash_attn.models.gpt import GPTLMHeadModel, remap_state_dict_hf_gpt2

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.distributed import all_gather_raw

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

torch.cuda API usage

torch.cuda.set_device(device)

template:47

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:49

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

tests/models/test_gpt_neox.py· 4

FlashAttention dependency

from flash_attn.models.gpt import GPTLMHeadModel

template:7

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.gpt_neox import gpt_neox_config_to_gpt2_config, remap_state_dict_hf_gpt_neox

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:42

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/models/test_gpt_parallel.py· 7

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:10

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

FlashAttention dependency

from flash_attn.losses.cross_entropy import CrossEntropyLoss

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.gpt import GPTLMHeadModel, shard_state_dict_tp

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.distributed import allreduce_sequence_parallel_grad

template:14

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:17

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:17

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=dtype):

template:112

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/models/test_gpt.py· 19

FlashAttention dependency

from flash_attn.models.gpt import (

template:6

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.generation import InferenceParams

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

.cuda() tensor/module move

model = model.cuda().to(dtype=dtype)

template:41

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

.cuda() tensor/module move

model_ref = GPT2LMHeadModelHF.from_pretrained(model_name).cuda()

template:43

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

.cuda() tensor/module move

model_hf = GPT2LMHeadModelHF.from_pretrained(model_name).cuda().to(dtype=dtype)

template:44

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

seqlens = torch.randint(max_seqlen // 2, max_seqlen + 1, (batch_size,), device="cuda")

template:53

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

0, config.vocab_size, (batch_size, max_seqlen), dtype=torch.long, device="cuda"

template:55

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

.cuda() tensor/module move

model = model.cuda().to(dtype=dtype)

template:98

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

.cuda() tensor/module move

model_ref = GPT2LMHeadModelHF.from_pretrained(model_name).cuda()

template:100

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

.cuda() tensor/module move

model_hf = GPT2LMHeadModelHF.from_pretrained(model_name).cuda().to(dtype=dtype)

template:101

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

seqlens = torch.randint(max_seqlen // 2, max_seqlen + 1, (batch_size,), device="cuda")

template:110

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

0, vocab_size_og, (batch_size, max_seqlen), dtype=torch.long, device="cuda"

template:112

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:148

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:285

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:348

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:395

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

FlashAttention dependency

from flash_attn.utils.generation import decode_speculative

template:424

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

model = GPTLMHeadModel(config, device="cuda", dtype=torch.float16)

template:466

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/models/test_gptj.py· 12

FlashAttention dependency

from flash_attn.models.gpt import GPTLMHeadModel

template:7

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.gptj import gptj_config_to_gpt2_config, remap_state_dict_hf_gptj

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.generation import update_graph_cache

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:33

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:93

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:118

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:123

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:137

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:148

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:155

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:166

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/models/test_llama.py· 19

FlashAttention dependency

from flash_attn.models.gpt import GPTLMHeadModel, combine_state_dicts_tp, shard_state_dict_tp

template:19

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.llama import (

template:20

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.distributed import all_gather_raw

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.generation import update_graph_cache

template:29

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:30

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:105

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:191

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

Device string "cuda"

device = "cuda"

template:295

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:320

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:325

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:346

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:357

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:364

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:375

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:407

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

torch.cuda API usage

torch.cuda.set_device(device)

template:448

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:501

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:510

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:538

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

tests/models/test_opt.py· 16

FlashAttention dependency

from flash_attn.models.gpt import GPTLMHeadModel

template:7

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.models.opt import opt_config_to_gpt2_config, remap_state_dict_hf_opt

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.generation import update_graph_cache

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.pretrained import state_dict_from_pretrained

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:39

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

seqlens = torch.randint(max_seqlen // 2, max_seqlen + 1, (batch_size,), device="cuda")

template:61

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

0, config.vocab_size, (batch_size, max_seqlen), dtype=torch.long, device="cuda"

template:63

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:108

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.synchronize()

template:151

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:161

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:171

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:181

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:192

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:197

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:204

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:209

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/models/test_vit.py· 2

FlashAttention dependency

from flash_attn.models.vit import vit_base_patch16_224 as flash_vit_base_patch16_224

template:5

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:19

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/modules/test_block_parallel.py· 7

NVIDIA Apex dependency

from apex.transformer import parallel_state, tensor_parallel

template:11

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

FlashAttention dependency

from flash_attn.modules.block import Block

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.mha import MHA, ParallelMHA

template:14

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.modules.mlp import FusedMLP, ParallelFusedMLP

template:15

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.utils.distributed import allreduce_sequence_parallel_grad

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:18

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:18

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/modules/test_embedding_parallel.py· 4

NVIDIA Apex dependency

from apex.transformer import parallel_state

template:8

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

FlashAttention dependency

from flash_attn.modules.embedding import GPT2Embeddings, ParallelGPT2Embeddings

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:12

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:12

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/modules/test_mha_parallel.py· 4

NVIDIA Apex dependency

from apex.transformer import parallel_state, tensor_parallel

template:9

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

FlashAttention dependency

from flash_attn.modules.mha import MHA, ParallelMHA

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:13

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:13

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/modules/test_mlp_parallel.py· 4

NVIDIA Apex dependency

from apex.transformer import parallel_state, tensor_parallel

template:7

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

FlashAttention dependency

from flash_attn.modules.mlp import GatedMlp, ParallelGatedMlp

template:9

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:11

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:11

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/ops/test_dropout_layer_norm.py· 15

FlashAttention dependency

from flash_attn.ops.layer_norm import (

template:7

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.ops.rms_norm import (

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

NVIDIA Apex dependency

from apex.normalization import FusedRMSNorm

template:21

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

NVIDIA Apex dependency

from apex.normalization.fused_layer_norm import fused_rms_norm_affine

template:22

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:27

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:27

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:70

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:180

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:257

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:374

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:440

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:597

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:782

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:991

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1165

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/ops/test_fused_dense_parallel.py· 4

NVIDIA Apex dependency

from apex.transformer import parallel_state, tensor_parallel

template:9

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

FlashAttention dependency

from flash_attn.ops.fused_dense import ColumnParallelLinear, FusedDense, FusedMLP, ParallelFusedMLP

template:10

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:12

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:12

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

tests/ops/test_fused_dense.py· 3

FlashAttention dependency

from flash_attn.ops.fused_dense import FusedDense, FusedMLP

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:17

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:103

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/ops/triton/test_layer_norm.py· 7

FlashAttention dependency

from flash_attn.ops.triton.layer_norm import (

template:8

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:16

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] >= 8

template:16

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:63

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:266

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=input_dtype):

template:324

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=input_dtype):

template:340

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/test_flash_attn_ck.py· 24

FlashAttention dependency

from flash_attn import (

template:7

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb

template:28

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:92

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:190

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:323

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:540

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:800

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:902

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1092

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1328

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1385

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn([batch_size, seqlen, nheads, d], dtype=dtype, device="cuda") * 5

template:1390

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn([batch_size, seqlen, nheads, d], dtype=dtype, device="cuda") * 3

template:1392

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1440

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn([batch_size, seqlen, nheads, d], dtype=dtype, device="cuda", requires_grad=True)

template:1446

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

g = torch.randn(seqlen, 2 * batch_size, nheads, d, dtype=dtype, device="cuda")[:, ::2]

template:1451

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1493

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1528

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1586

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1586

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:1591

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1637

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1637

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:1642

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/test_flash_attn_triton_amd.py· 47

FlashAttention dependency

from flash_attn import (

template:7

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.bert_padding import pad_input, unpad_input

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.flash_attn_interface import _get_block_size_n

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm75 = torch.cuda.get_device_capability("cuda") == (7, 5)

template:36

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm75 = torch.cuda.get_device_capability("cuda") == (7, 5)

template:36

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] == 8

template:37

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] == 8

template:37

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

is_sm80 = torch.cuda.get_device_capability("cuda") == (8, 0)

template:38

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm80 = torch.cuda.get_device_capability("cuda") == (8, 0)

template:38

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

is_sm90 = torch.cuda.get_device_capability("cuda") == (9, 0)

template:39

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm90 = torch.cuda.get_device_capability("cuda") == (9, 0)

template:39

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

mask = torch.rand(nrow, ncol, device="cuda") < sparsity

template:393

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

if seqlen >= 2048 and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30:

template:602

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if seqlen >= 2048 and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30:

template:602

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:604

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

if seqlen >= 2048 and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30:

template:751

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if seqlen >= 2048 and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30:

template:751

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:753

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:927

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:927

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:932

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1199

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1199

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:1204

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1511

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1511

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:1516

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1624

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1624

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:1629

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1800

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1963

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2231

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2283

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn([batch_size, seqlen, nheads, d], dtype=dtype, device="cuda") * 5

template:2288

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn([batch_size, seqlen, nheads, d], dtype=dtype, device="cuda") * 3

template:2290

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2340

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn([batch_size, seqlen, nheads, d], dtype=dtype, device="cuda", requires_grad=True)

template:2346

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

g = torch.randn(seqlen, 2 * batch_size, nheads, d, dtype=dtype, device="cuda")[:, ::2]

template:2351

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2393

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:2451

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:2451

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:2456

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:2510

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:2510

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:2515

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/test_flash_attn.py· 49

FlashAttention dependency

from flash_attn import (

template:7

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.bert_padding import pad_input, unpad_input

template:16

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.flash_attn_interface import _get_block_size_n

template:17

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb

template:18

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm75 = torch.cuda.get_device_capability("cuda") == (7, 5)

template:23

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm75 = torch.cuda.get_device_capability("cuda") == (7, 5)

template:23

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda")[0] == 8

template:24

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda")[0] == 8

template:24

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

is_sm80 = torch.cuda.get_device_capability("cuda") == (8, 0)

template:25

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm80 = torch.cuda.get_device_capability("cuda") == (8, 0)

template:25

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

is_sm90 = torch.cuda.get_device_capability("cuda") == (9, 0)

template:26

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm90 = torch.cuda.get_device_capability("cuda") == (9, 0)

template:26

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

mask = torch.rand(nrow, ncol, device="cuda") < sparsity

template:378

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

if seqlen >= 2048 and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30:

template:587

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if seqlen >= 2048 and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30:

template:587

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:589

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

if seqlen >= 2048 and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30:

template:736

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

if seqlen >= 2048 and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30:

template:736

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:738

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:908

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:908

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:913

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1177

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1177

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:1182

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1485

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1485

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:1490

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1598

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:1598

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:1603

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1770

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:1933

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2200

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2251

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

q = torch.randn([batch_size, seqlen, nheads, d], dtype=dtype, device="cuda") * 5

template:2256

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn([batch_size, seqlen, nheads, d], dtype=dtype, device="cuda") * 3

template:2258

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2307

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

torch.randn([batch_size, seqlen, nheads, d], dtype=dtype, device="cuda", requires_grad=True)

template:2313

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

g = torch.randn(seqlen, 2 * batch_size, nheads, d, dtype=dtype, device="cuda")[:, ::2]

template:2318

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:2359

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:2416

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:2416

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:2421

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:2474

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

and torch.cuda.get_device_properties("cuda").total_memory <= 16 * 2**30

template:2474

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:2479

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

FlashAttention dependency

from flash_attn.flash_attn_interface import _flash_attn_varlen_forward

template:2531

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

device = "cuda"

template:2533

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

tests/test_rotary.py· 13

Triton dependency

import triton

template:9

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb, apply_rotary_emb_torch

template:11

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.layers.rotary import apply_rotary_emb_qkv_, apply_rotary_emb_kv_

template:12

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

FlashAttention dependency

from flash_attn.bert_padding import pad_input, unpad_input

template:13

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

Device string "cuda"

is_sm8x = torch.cuda.get_device_capability("cuda") >= (8, 0)

template:15

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

is_sm8x = torch.cuda.get_device_capability("cuda") >= (8, 0)

template:15

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

device = "cuda"

template:66

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:121

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:187

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:235

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

device = "cuda"

template:284

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Triton dependency

from triton.runtime.jit import JITFunction

template:288

OpenAI Triton has an AMD backend. Kernels usually run on ROCm with a matching Triton build; occasional tuning is needed.

RecommendInstall the ROCm-enabled Triton build.

FlashAttention dependency

from flash_attn.ops.triton.rotary import rotary_kernel

template:289

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

tests/test_util.py· 1

FlashAttention dependency

from flash_attn.bert_padding import pad_input, unpad_input

template:5

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

training/src/callbacks/norm_monitor.py· 1

NVIDIA Apex dependency

from apex.contrib.layer_norm import FastLayerNorm

template:17

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

training/src/metrics/perplexity.py· 1

FlashAttention dependency

from flash_attn.losses.cross_entropy import CrossEntropyLoss

template:14

FlashAttention has an official ROCm/CK implementation. Swapping the wheel (or using PyTorch SDPA) is mechanical.

RecommendInstall the ROCm FlashAttention build or fall back to torch SDPA.

training/src/optim/param_grouping.py· 1

NVIDIA Apex dependency

from apex.contrib.layer_norm import FastLayerNorm

template:8

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

training/src/utils/ddp_zero2.py· 3

NVIDIA Apex dependency

from apex.contrib.optimizers.distributed_fused_adam import DistributedFusedAdam

template:11

Apex fused optimizers/AMP are CUDA-specific. ROCm PyTorch ships native AMP (torch.cuda.amp / torch.amp) covering most uses.

RecommendReplace Apex AMP with native torch.amp; use ROCm apex for fused ops.

torch.cuda API usage

torch.cuda.empty_cache()

template:134

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

Device string "cuda"

map_location='cuda'

template:143

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

training/src/utils/gpu_affinity.py· 1

pynvml (NVML) dependency

import pynvml

template:7

pynvml queries NVIDIA-only NVML. On AMD the equivalent is amdsmi / rocm-smi; the query calls map across mechanically.

RecommendReplace pynvml calls with amdsmi (ROCm SMI Python bindings).

benchmarks/tune_ex2_emu.py· 5

torch.cuda API usage

major, minor = torch.cuda.get_device_capability()

template:84

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

print(f"GPU: {torch.cuda.get_device_name()}, SM{sm}, is_sm103={is_sm103}")

template:87

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

idx = torch.cuda.current_device() if torch.cuda.is_available() else 0

template:100

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

if torch.cuda.is_available():

template:102

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

return str(torch.cuda.current_device())

template:103

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

csrc/layer_norm/setup.py· 1

torch.cuda API usage

if not torch.cuda.is_available():

template:61

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/cute/bench_utils.py· 2

Device string "cuda"

row_idx = torch.arange(seqlen_q, device="cuda")

template:32

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

workspace = torch.empty(graph.get_workspace_size(), device="cuda", dtype=torch.uint8)

template:142

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

flash_attn/cute/benchmark.py· 13

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:16

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:44

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:86

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:100

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:215

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:225

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:245

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:259

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.reset_peak_memory_stats()

template:260

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:261

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:263

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

mem = torch.cuda.max_memory_allocated() / ((2**20) * 1000)

template:264

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.empty_cache()

template:267

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/cute/testing.py· 1

Device string "cuda"

(batch, seqlen_q, seqlen_k), False, device="cuda"

template:423

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

flash_attn/flash_blocksparse_attn_interface.py· 8

torch.cuda API usage

rng_state = torch.cuda.get_rng_state() if dropout_p > 0 else None

template:90

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

cur_rng_state = torch.cuda.get_rng_state()

template:114

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.set_rng_state(rng_state)

template:115

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.set_rng_state(cur_rng_state)

template:131

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

rng_state = torch.cuda.get_rng_state() if dropout_p > 0 else None

template:141

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

cur_rng_state = torch.cuda.get_rng_state()

template:165

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.set_rng_state(rng_state)

template:166

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.set_rng_state(cur_rng_state)

template:181

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/ops/activations.py· 2

torch.cuda API usage

swiglu_fwd = torch.cuda.jiterator._create_jit_fn(swiglu_fwd_codestring)

template:119

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

swiglu_bwd = torch.cuda.jiterator._create_multi_output_jit_fn(swiglu_bwd_codestring, num_outputs=2)

template:120

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/utils/benchmark.py· 13

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:16

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:44

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:86

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:100

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:215

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:225

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

with torch.autocast(device_type="cuda", dtype=amp_dtype, enabled=amp):

template:245

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

torch.cuda.empty_cache()

template:259

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.reset_peak_memory_stats()

template:260

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:261

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:263

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

mem = torch.cuda.max_memory_allocated() / ((2**20) * 1000)

template:264

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.empty_cache()

template:267

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/utils/generation.py· 12

torch.cuda API usage

start = torch.cuda.Event(enable_timing=enable_timing)

template:187

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

end = torch.cuda.Event(enable_timing=enable_timing)

template:188

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:203

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:415

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.synchronize()

template:551

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

cache.mempool = torch.cuda.graphs.graph_pool_handle()

template:671

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

s = torch.cuda.Stream()

template:704

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

s.wait_stream(torch.cuda.current_stream())

template:705

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

with torch.cuda.stream(s):

template:706

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

torch.cuda.current_stream().wait_stream(s)

template:720

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

graph = torch.cuda.CUDAGraph()

template:723

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

torch.cuda API usage

with torch.cuda.graph(graph, pool=mempool):

template:724

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

flash_attn/utils/torch.py· 1

Device string "cuda"

kwargs["device_type"] = "cuda"

template:8

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

hopper/flash_attn_interface.py· 3

Device string "cuda"

@torch.library.custom_op("flash_attn_3::_flash_attn_forward", mutates_args=(), device_types="cuda")

template:59

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

Device string "cuda"

@torch.library.custom_op("flash_attn_3::_flash_attn_backward", mutates_args=("dq", "dk", "dv"), device_types="cuda")

template:258

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.

torch.cuda API usage

cap = torch.cuda.get_device_capability(q.device)

template:369

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

hopper/test_torch_compile_and_export.py· 2

.cuda() tensor/module move

model = EfficienctMultiHeadAttention(embedding_dim, num_heads).cuda().bfloat16()

template:40

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

.cuda() tensor/module move

input_tensor = torch.randn(batch_size, sequence_length, embedding_dim).cuda().bfloat16()

template:41

.cuda() on tensors and modules is honoured by ROCm PyTorch and moves data to the AMD GPU. No change needed.

RecommendNo change required on ROCm PyTorch.

training/src/callbacks/gpu_affinity.py· 1

torch.cuda API usage

nproc_per_node = torch.cuda.device_count()

template:24

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

training/src/utils/ddp_zero1.py· 1

torch.cuda API usage

torch.cuda.empty_cache()

template:97

Calls under torch.cuda.* are aliased by ROCm builds of PyTorch. The same code runs on AMD Instinct GPUs unchanged — torch.cuda.is_available(), streams, events and AMP all map onto HIP.

RecommendNo change required. Install the ROCm build of PyTorch (pip install torch --index-url https://download.pytorch.org/whl/rocm6.1).

training/src/utils/distributed.py· 1

Device string "cuda"

device = torch.device('cuda')

template:88

The literal device string "cuda" resolves to the active AMD GPU on ROCm PyTorch. .to("cuda") / device="cuda" need no edits.

RecommendNo change required on ROCm PyTorch.