Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
52 commits
Select commit Hold shift + click to select a range
7bfd03f
feat(stinkytofu): add gfx1250 replay hazard protection (#10104)
cycheng Aug 7, 2026
8ee23cc
fix(rocalution): add missing break in complex32 case of rocsparseio r…
adityasingh2400 Aug 7, 2026
ed4cc48
build(deps): bump cryptography from 48.0.1 to 50.0.0 in /projects/roc…
dependabot[bot] Aug 7, 2026
c6c1b67
docs(rocsparse): Fix docs build failure caused by missing define (#10…
jsandham Aug 7, 2026
10b7238
perf(rocprim): remove one extra scan from nth-element (#8306)
Saiyang-Zhang Aug 7, 2026
5bc2e15
feat(hipdnn): implement remaining cuDNN frontend graph nodes (#9738)
mousdahl-amd Aug 7, 2026
43b8dca
perf(hipblaslt) Update 44 f32 tn sizes for gfx942 80cu grid (#9160)
m-kim Aug 7, 2026
26c96a3
ci(hip-kernel-provider): stop installing rocKE Python tests in provid…
BrianHarrisonAMD Aug 7, 2026
a9ade2b
feat(hipblaslt-provider): support OCP MX FP4 (E2M1) dequantize + GEMM…
averinevg Aug 7, 2026
84bb57b
build(deps): bump rocm-docs-core from 1.32.1 to 1.38.0 in /projects/h…
dependabot[bot] Aug 7, 2026
49d85ee
build(deps): bump gitpython from 3.1.57 to 3.1.58 in /shared/tensile/…
dependabot[bot] Aug 7, 2026
311eb7c
perf(rocsparse): fit bsrmv/bsrxmv general-path wave32 geometry to blo…
kliegeois Aug 7, 2026
54f224d
fix(rocsparse): Fix SpMM batching with COO segmented algorithm when b…
jsandham Aug 7, 2026
cf18605
feat(hipcub, rocprim, rocthrust) Revert CCCL 3.0 changes (#10464)
stanleytsang-amd Aug 7, 2026
16ad891
Revert "perf(hipblaslt) Update 44 f32 tn sizes for gfx942 80cu grid" …
m-kim Aug 7, 2026
8244e10
fix[stinkytofu] Fix co-execution V_NOP handling and add DAG scheduler…
darrenhsieh-amd Aug 7, 2026
1ef3250
Add Rocke platform PR label (#10372)
JohnUzoka Aug 7, 2026
62657e8
feat: build FFT clients with amdgcnspirv arch
evetsso Aug 7, 2026
7e90708
feat(ci): enable automatic stage reuse for TheRock builds (#10533)
geomin12 Aug 7, 2026
8d24cb4
feat(complex): add gfx1250 CGEMM 'TDM' library logic (NN/NT/TT) (#10509)
mengzcai Aug 7, 2026
795aae5
fix(hipfft-test): initialize and compute reference results only if re…
regan-amd Aug 8, 2026
01f686a
test(hipdnn): fix ragged tensor int32-narrowing test failing in debug…
CMiservaAMD Aug 8, 2026
ec65811
feat(rocke): conv pointwise optimization (#10115)
bartekxk Aug 8, 2026
a2d4fa7
test(rocke): bench cases conv (#10059)
bartekxk Aug 8, 2026
e3bdcef
docs(rpp): refresh docs for monorepo move, TheRock delivery, and 3.2.…
kiritigowda Aug 8, 2026
f61ae65
ci(hipblaslt): support different OpenMP locations (AIHPBLAS-3605) (#1…
newling Aug 9, 2026
1f2fd7b
feat(hipsparse): add generic SpGEAM API for sparse matrix-matrix addi…
kliegeois Aug 10, 2026
535da60
fix(MIOpen): correct deterministic SetNextValue polarity for grouped …
JonathanLichtnerAMD Aug 10, 2026
1289f6c
fix(theRock artifacts): Exclude optional LFS matrices from hipthreads…
marsavicc Aug 10, 2026
2531c64
fix(stinkytofu): guard clang-tidy setup behind ENABLE_CLANG_TIDY (#10…
KKyang Aug 10, 2026
db70f7d
feat(stinkytofu): Add RemoveDscntPass for WMMA Loops (#10389)
hcman2 Aug 10, 2026
5f98f20
fix(rocisa): rebuild on import and correct staleness false positive (…
KKyang Aug 10, 2026
1e27aa6
test(hipthreads): prefer ROCm bin DLLs for lit tests on Windows (#101…
marsavicc Aug 10, 2026
6c627c2
feat(ck) [CK] Wavelet gemm pipeline for conv fwd (#10179)
jakpiase Aug 10, 2026
c0d4441
refactor(stinkytofu): unify tie-break candidate selection in CDNA5 sc…
KKyang Aug 10, 2026
6a7c678
doc(rocBLAS) : adds testing guide (#10502)
TorreZuk Aug 10, 2026
cb9f216
fix(primbench): add fallback clock rate calculation (#10194)
NguyenNhuDi Aug 10, 2026
76ab1b6
feat(rocBLAS/Tensile): new version bumps (#10369)
NaveenElumalaiAMD Aug 10, 2026
966f7b8
[hipblaslt] Equality tuning for BF16 TN (gfx1250)
minsukim-amd Aug 10, 2026
7730ced
Updated gfx1250 Equality BBS_TN liblogic yaml file
minsukim-amd Aug 10, 2026
1c1b5c3
docs(hipsparse): fix documentation warnings (#10578)
jsandham Aug 10, 2026
a315b98
fix(hiptensor): Fix bin installation paths for tests on Windows (#10491)
evedovelli Aug 10, 2026
2cd2532
Merge branch 'develop' into users/minsukim/gfx1250_equality_bbstn_202…
minsukim-amd Aug 10, 2026
15f8fb3
Removed 2 invalid kernels
minsukim-amd Aug 10, 2026
c4a1de3
feat(hipdnn): declare HIPDNN_ENABLE_KERNEL_INGESTOR option (#10576)
SamuelReeder Aug 10, 2026
c4c1368
gfx1250 Equality tuning for BBS NN, NT
aferoz21 Aug 10, 2026
fcb9efa
fix(ci): add actions:read permission to multi-arch workflows (#10573)
geomin12 Aug 10, 2026
4b253ee
Fix dict-format serialization of gfx1250 BBS NN/NT Equality logic
minsukim-amd Aug 10, 2026
3a08f1b
[hipblaslt] Equality tuning for BF16, MXF8, MXF4, F8 MAF (#10250)
minsukim-amd Aug 10, 2026
535d458
fix(tensilelite): Fix LocalRead comments to list all LDS mem tokens (…
hcman2 Aug 11, 2026
85b84bb
Remove invalid kernels of gfx1250 BBS NN and BBS NT type
aferoz21 Aug 11, 2026
bab3162
Merge remote-tracking branch 'origin/develop' into users/minsukim/gfx…
minsukim-amd Aug 11, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
7 changes: 7 additions & 0 deletions .github/labeler.yml
Original file line number Diff line number Diff line change
Expand Up @@ -24,13 +24,20 @@
- changed-files:
- any-glob-to-any-file: 'dnn-providers/hip-kernel-provider/**/*'

# These labels help distinguish which subteams are responsible for reviews for changes made in rocKE

"rocke-lib: attention":
- changed-files:
- any-glob-to-any-file: 'dnn-providers/hip-kernel-provider/rocke/library/**/*attention*'
- any-glob-to-any-file: 'dnn-providers/hip-kernel-provider/rocke/library/**/*fmha*'
- any-glob-to-any-file: 'dnn-providers/hip-kernel-provider/rocke/library/**/*sdpa*'
- any-glob-to-any-file: 'dnn-providers/hip-kernel-provider/rocke/library/builders/**/attention/**/*'

# This label serves to distiguish work being done in RocKE platform, for the reviewing team and internal tools
"rocke: platform":
- changed-files:
- any-glob-to-any-file: 'dnn-providers/hip-kernel-provider/rocke/platform/**/*'

"project: hipdnn-integration-tests":
- changed-files:
- any-glob-to-any-file: 'dnn-providers/integration-tests/**/*'
Expand Down
36 changes: 31 additions & 5 deletions .github/scripts/run_rocjitsu_hipblaslt_race_check.sh
Original file line number Diff line number Diff line change
Expand Up @@ -139,13 +139,38 @@ if [[ ! -x "${TENSILELITE_CLIENT}" ]]; then
exit 1
fi

# The fetched compiler artifact installs the host OpenMP runtime here.
LLVM_HOST_RUNTIME_DIR="${ROCM_PATH}/lib/llvm/lib/x86_64-unknown-linux-gnu"
if [[ ! -f "${LLVM_HOST_RUNTIME_DIR}/libomp.so" ]]; then
echo "OpenMP runtime not found: ${LLVM_HOST_RUNTIME_DIR}/libomp.so" >&2
# The compiler artifact layout depends on whether TheRock enables LLVM's host
# per-target runtime directories. Prefer the flat lib/llvm/lib layout, then the
# matching host-triple directory when per-target runtime directories are used.
LLVM_RUNTIME_ROOT="${ROCM_PATH}/lib/llvm/lib"
if [[ ! -d "${LLVM_RUNTIME_ROOT}" ]]; then
echo "LLVM runtime directory not found: ${LLVM_RUNTIME_ROOT}" >&2
exit 1
fi

LIBOMP_CANDIDATES=(
"${LLVM_RUNTIME_ROOT}/libomp.so"
"${LLVM_RUNTIME_ROOT}/$(uname -m)-unknown-linux-gnu/libomp.so"
)
LIBOMP_PATH=""
for candidate in "${LIBOMP_CANDIDATES[@]}"; do
if [[ -e "${candidate}" ]]; then
LIBOMP_PATH="${candidate}"
break
fi
done

if [[ -z "${LIBOMP_PATH}" ]]; then
echo "OpenMP runtime not found at either expected path:" >&2
printf " %s\n" "${LIBOMP_CANDIDATES[@]}" >&2
exit 1
fi
LLVM_HOST_RUNTIME_DIR="$(dirname "${LIBOMP_PATH}")"
LLVM_RUNTIME_LIBRARY_PATH="${LLVM_RUNTIME_ROOT}"
if [[ "${LLVM_HOST_RUNTIME_DIR}" != "${LLVM_RUNTIME_ROOT}" ]]; then
LLVM_RUNTIME_LIBRARY_PATH="${LLVM_RUNTIME_LIBRARY_PATH}:${LLVM_HOST_RUNTIME_DIR}"
fi

select_rocjitsu_target

if [[ -z "${ROCJITSU_CONFIG}" || ! -f "${ROCJITSU_CONFIG}" ]]; then
Expand All @@ -161,7 +186,7 @@ mkdir -p "${RACE_REPORT_DIR}"
# that use it to find the ROCm install root.
export ROCM_PATH
export PATH="${ROCM_PATH}/bin:${ROCM_PATH}/lib/llvm/bin:${PATH}"
export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/rocm_sysdeps/lib:${ROCM_PATH}/lib/llvm/lib:${LLVM_HOST_RUNTIME_DIR}:${LD_LIBRARY_PATH:-}"
export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/lib/rocm_sysdeps/lib:${LLVM_RUNTIME_LIBRARY_PATH}:${LD_LIBRARY_PATH:-}"
export PYTHONPATH="${TENSILELITE_ROOT}${PYTHONPATH:+:${PYTHONPATH}}"

echo "ROCM_PATH=${ROCM_PATH}"
Expand All @@ -174,6 +199,7 @@ echo "HIPBLASLT_BENCH=${HIPBLASLT_BENCH}"
echo "TENSILELITE_ROOT=${TENSILELITE_ROOT}"
echo "TENSILE_DRIVER=${TENSILE_DRIVER}"
echo "TENSILELITE_CLIENT=${TENSILELITE_CLIENT}"
echo "LIBOMP_PATH=${LIBOMP_PATH}"
echo "RACE_REPORT_DIR=${RACE_REPORT_DIR}"
echo "PATH=${PATH}"
echo "LD_LIBRARY_PATH=${LD_LIBRARY_PATH}"
Expand Down
2 changes: 2 additions & 0 deletions .github/workflows/hipdnn-superbuild-ci.yml
Original file line number Diff line number Diff line change
Expand Up @@ -120,6 +120,7 @@ jobs:
-DROCM_LIBS_ENABLE_ROOT_CTEST=ON \
-DENABLE_CLANG_FORMAT=OFF \
-DHIPKERNELPROVIDER_ENABLE_ROCKE=ON \
-DHIPDNN_ENABLE_CUDNN_COMPATIBILITY=ON \
-DENABLE_CLANG_TIDY=OFF

- name: Build
Expand Down Expand Up @@ -280,6 +281,7 @@ jobs:
-DENABLE_CLANG_FORMAT=OFF \
-DENABLE_CLANG_TIDY=ON \
-DHIPKERNELPROVIDER_ENABLE_ROCKE=ON \
-DHIPDNN_ENABLE_CUDNN_COMPATIBILITY=ON \
-DCMAKE_CXX_CLANG_TIDY=/usr/bin/clang-tidy-20

- name: Build
Expand Down
1 change: 1 addition & 0 deletions .github/workflows/therock-multi-arch-ci-asan-nightly.yml
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,7 @@ on:

permissions:
contents: read
actions: read

concurrency:
# For scheduled runs, cancel any in-progress nightly so only one runs at a time.
Expand Down
1 change: 1 addition & 0 deletions .github/workflows/therock-multi-arch-ci-asan.yml
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,7 @@ on:

permissions:
contents: read
actions: read

concurrency:
# Only allow one run at a time for 'pull_request' events.
Expand Down
1 change: 1 addition & 0 deletions .github/workflows/therock-multi-arch-ci-nightly.yml
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,7 @@ on:

permissions:
contents: read
actions: read

concurrency:
# A PR number if a pull request and otherwise the commit hash. This cancels
Expand Down
14 changes: 8 additions & 6 deletions .github/workflows/therock-multi-arch-ci.yml
Original file line number Diff line number Diff line change
Expand Up @@ -64,6 +64,7 @@ on:

permissions:
contents: read
actions: read

concurrency:
# A PR number if a pull request and otherwise the commit hash. This cancels
Expand Down Expand Up @@ -148,16 +149,17 @@ jobs:
windows_amdgpu_families: ${{ inputs.windows_amdgpu_families || 'gfx110X' }}
linux_test_labels: ${{ inputs.linux_test_labels || '' }}
windows_test_labels: ${{ inputs.windows_test_labels || '' }}
# HARDCODED: Copy all stages except math-libs from TheRock baseline
# rocm-libraries only affects math-libs and ml-libs source sets
# Note: compiler-runtime excluded due to hardcoded CMake paths (fftw3, etc.)
# TODO(geomin12): Enable automatic stage reuse once dry-run mode issues resolved
prebuilt_stages: ${{ inputs.prebuilt_stages || 'runtime-tests,comm-libs,storage-libs,debug-tools,dctools-core,profiler-apps,media-libs,wsl-rocdxg' }}
baseline_run_id: ${{ inputs.baseline_run_id || '30989762651' }}
# Manual override for prebuilt stages (empty = use automatic stage reuse)
prebuilt_stages: ${{ inputs.prebuilt_stages || '' }}
baseline_run_id: ${{ inputs.baseline_run_id || '' }}
# Enable automatic stage reuse - TheRock finds commit-compatible baselines
# and determines which stages to rebuild based on changed files
stage_reuse_mode: reuse-stage
repository: ROCm/TheRock
# Concrete SHA resolved once by resolve_therock_ref; setup re-exports the
# checked-out commit as needs.setup.outputs.ref for downstream jobs.
ref: ${{ needs.resolve_therock_ref.outputs.therock_ref }}
# Pass external repo info for stage-impact analysis
external_repo: '{"repository":"${{ github.repository }}","ref":"${{ github.sha }}"}'

linux_build_and_test:
Expand Down
23 changes: 23 additions & 0 deletions dnn-providers/hip-kernel-provider/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,7 @@ list(APPEND CMAKE_MODULE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/cmake

option(HIPKERNELPROVIDER_ENABLE_TESTS "Build tests" ON)
option(HIPKERNELPROVIDER_ENABLE_ROCKE "Build rocKE engine and smoke tests" OFF)
option(HIPDNN_ENABLE_KERNEL_INGESTOR "Enable the generic kernel ingestor" OFF)
option(ENABLE_ASM_SDPA_ENGINE "Build the ASM SDPA engine" ON)
option(ENABLE_HIP_MLOPS_ENGINE "Build the HIP kernel engine" ON)

Expand Down Expand Up @@ -194,6 +195,28 @@ if(HIPKERNELPROVIDER_ENABLE_ROCKE)
set(ROCKE_INSTALL_PROVIDER_CTEST_SUBDIR "hip_kernel_provider")
set(ROCKE_INSTALL_STANDALONE_CTEST OFF)

# TEMPORARY: do not install the rocKE Python CI tests (rocke_python_import_smoke,
# rocke_golden_static, rocke_pytest) in provider mode.
#
# These three CTest entries import the `rocke` package from the test artifact.
# ROCm/TheRock#6849 removed the "bin/hip_kernel_provider/rocke/**" artifact
# include on the assumption that the staged rocKE wheels replaced it, but the
# wheels are never staged (see below), so neither import source survives and
# every rocKE Python entry fails with `ModuleNotFoundError: No module named
# 'rocke'`. rocm-libraries builds this component with
# HIPKERNELPROVIDER_ENABLE_ROCKE=ON, so the breakage lands on every
# hipkernelprovider PR here while TheRock's own CI (flag defaults OFF) stays green.
#
# This matches what the ROCKE_INSTALL_PYTHON_TESTS option already documents for
# provider mode; nothing was actually setting it. The C++ rocKE smoke tests are
# unaffected and keep running.
#
# Revert this line once ROCm/TheRock#7145 has landed and the wheel-staging order
# bug in rocke/CMakeLists.txt is fixed (ROCKE_BUILD_PYENV / ROCKE_WHEEL_DIR are
# declared after add_subdirectory(platform), so the wheel install() never
# registers). Tracked by ROCm/rocm-libraries#10497.
set(ROCKE_INSTALL_PYTHON_TESTS OFF)

add_subdirectory(rocke)
else()
message(STATUS "rocKE features disabled")
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -177,6 +177,32 @@ int rocke_conv_problem_k_gemm(const rocke_conv_problem_t* p);
* int never hits. */
long long rocke_conv_problem_flops(const rocke_conv_problem_t* p);

/* ConvProblem.is_pointwise property:
* Y == 1 and X == 1 and sH == 1 and sW == 1 and pH == 0 and pW == 0
* and (not is_3d or (Z == 1 and sD == 1 and pD == 0))
* For pointwise convolutions A, B, D are flat 2-D matrices and the
* coordinate-transform DAG reduces to multiply+add. */
static inline bool rocke_conv_problem_is_pointwise(const rocke_conv_problem_t* p)
{
if(p->Y != 1 || p->X != 1 || p->sH != 1 || p->sW != 1 || p->pH != 0 || p->pW != 0)
return false;
if(p->is_3d && (p->Z != 1 || p->sD != 1 || p->pD != 0))
return false;
return true;
}

/* ConvProblem.cpg property: C / groups (channels per group). */
static inline int rocke_conv_problem_cpg(const rocke_conv_problem_t* p)
{
return p->groups > 1 ? p->C / p->groups : p->C;
}

/* ConvProblem.kpg property: K / groups (output channels per group). */
static inline int rocke_conv_problem_kpg(const rocke_conv_problem_t* p)
{
return p->groups > 1 ? p->K / p->groups : p->K;
}

/* ConvProblem.short() ->
* 2-D: f"N{N}H{Hi}W{Wi}C{C}_K{K}Y{Y}X{X}"
* 3-D: f"N{N}D{Di}H{Hi}W{Wi}C{C}_K{K}Z{Z}Y{Y}X{X}"
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -225,13 +225,44 @@ typedef struct rocke_conv_build_ctx
int load_vec; /* _choose_load_vec(spec) */

/* ---- coordinate-transform descriptors ---- */
rocke_tensor_descriptor_t* A_desc; /* make_a_descriptor(p, decompose_m) */
rocke_tensor_descriptor_t* B_desc; /* make_b_descriptor(p) */
rocke_tensor_descriptor_t* A_desc; /* make_a_descriptor(p, decompose_m); NULL if pointwise */
rocke_tensor_descriptor_t* B_desc; /* make_b_descriptor(p); NULL if pointwise */
/* D_desc is built lazily inside the epilogue phase (Python builds it per
* epilogue fn); held here so an epilogue_override and the two stock
* epilogues share one instance. NULL until the epilogue phase populates. */
rocke_tensor_descriptor_t* D_desc;

/* ---- pointwise fast-path (Python p.is_pointwise) ---- *
* When true, A_desc / B_desc / D_desc are NULL and flat multiply+add
* arithmetic is used instead of the coordinate-transform DAG.
* Mirrors: if p.is_pointwise: A_desc = None; B_desc = None ...
*
* Forward conv: c_M_pw=p.M, c_C_pw=p.cpg, c_K_pw=p.kpg.
* Wgrad: c_M_pw=wg_M, c_C_pw=p.cpg, c_K_pw=p.kpg,
* c_wgK_pw=wg_K, c_wgN_pw=wg_N. */
bool is_pointwise;
int c_M_pw; /* forward: p.M; wgrad: wg_M */
int c_C_pw; /* forward: p.cpg; wgrad: p.cpg (input channels per group) */
int c_K_pw; /* forward: p.kpg; wgrad: p.kpg (output channels per group) */
int c_wgK_pw; /* wgrad only: wg_K (= N*Ho*Wo, reduction dim); 0 for forward */
int c_wgN_pw; /* wgrad only: wg_N (= Y*X*C, N dim); 0 for forward */

/* ---- pointwise IR constants (emitted before buffer resources, matching Python) ---- *
* Python emits _c_C_ir, _c_K_ir, _c_M_ir, _always_valid (in that order)
* before make_buffer_resource calls. We pre-emit and cache them so the
* descriptor callbacks reuse rather than re-emit them. */
rocke_value_t* ir_c_C_pw; /* const_i32(cpg) — first (fwd) / _c_K_ir=kpg (wgrad) */
rocke_value_t* ir_c_K_pw; /* const_i32(kpg) — second (fwd) / _c_C_ir=cpg (wgrad) */
rocke_value_t* ir_c_M_pw; /* const_i32(M/wg_M) — third (fwd) / _c_wgM_ir (wgrad) */
rocke_value_t* ir_always_valid; /* const_i32(1) — fourth (fwd) / _c_wgN_ir (wgrad) */
rocke_value_t* ir_c_wgN_pw; /* wgrad only: const_i32(wg_N) — fifth */

/* ---- optional B-descriptor override (wgrad only) ---- *
* When non-NULL, rocke_conv_emit_load_phase calls this instead of
* rocke_conv_b_descriptor. Used by wgrad to provide a different pointwise
* formula for the X (input) B-tile descriptor. */
rocke_loads_descriptor_fn b_descriptor_fn;

/* ---- buffer resources (CK-Tile views over A/B/D) ---- */
rocke_conv_buffer_resource_t a_buf_rsrc; /* make_buffer_resource(b, A, A_bytes) */
rocke_conv_buffer_resource_t b_buf_rsrc; /* make_buffer_resource(b, Bp, B_bytes)*/
Expand Down Expand Up @@ -436,11 +467,11 @@ void rocke_conv_emit_direct_epilogue(rocke_ir_builder_t* b,
rocke_value_t* const* accs,
int num_accs,
const rocke_warp_grid_t* grid,
rocke_value_t* d_rsrc);
rocke_value_t* d_rsrc,
rocke_value_t* ir_c_K_pw);

/* _emit_direct_epilogue_wmma(b, spec, op, accs[n], warp_m_idx, warp_n_idx, lane,
* block_m_off, block_n_off, d_rsrc, c0): WMMA per-lane fp16 store via the op's
* c_layout map + the D descriptor. */
* block_m_off, block_n_off, d_rsrc, c0, ir_c_K_pw): WMMA per-lane fp16 store. */
void rocke_conv_emit_direct_epilogue_wmma(rocke_ir_builder_t* b,
const rocke_implicit_gemm_conv_spec_t* spec,
const rocke_mmaop_t* op,
Expand All @@ -452,16 +483,18 @@ void rocke_conv_emit_direct_epilogue_wmma(rocke_ir_builder_t* b,
rocke_value_t* block_m_off,
rocke_value_t* block_n_off,
rocke_value_t* d_rsrc,
rocke_value_t* c0);
rocke_value_t* c0,
rocke_value_t* ir_c_K_pw);

/* _emit_cshuffle_epilogue(b, spec, accs[n], grid, d_rsrc): LDS-staged cshuffle
* store via CShuffleEpilogue.from_grid + the D descriptor addr_fn. */
/* _emit_cshuffle_epilogue(b, spec, accs[n], grid, d_rsrc, ir_c_K_pw): LDS-staged
* cshuffle store via CShuffleEpilogue.from_grid + the D descriptor addr_fn. */
void rocke_conv_emit_cshuffle_epilogue(rocke_ir_builder_t* b,
const rocke_implicit_gemm_conv_spec_t* spec,
rocke_value_t* const* accs,
int num_accs,
const rocke_warp_grid_t* grid,
rocke_value_t* d_rsrc);
rocke_value_t* d_rsrc,
rocke_value_t* ir_c_K_pw);

/* ----- driver-internal ctx population (the build prologue, lines 787-1032) ----
* Splitting the long prologue out of the public entry keeps the glue TU small;
Expand Down
Loading
Loading