Andreas Kieslinger
e947228222
* Adds initial PDL setup. * Adds PDL barriers based on simple heuristic: place "sync" before first input pointer access, and "launch" after last write, e.g. to tensors like dst. * Further optimization pass of the first half of kernels * Optimized PDL barriers for the second batch of kernels * Further refinements after rebase. * Moves pdl logic to separate function, removes some whitespace * Strips post-hoc PDL logic * Adds stream capture PDL setup. Enrolls quantize_q8_1 to leverage pdl to overlap execution with previous kernels * Enrolls mul_mat_vec_q, rms_norm_f32 and k_bin_bcast (partly) into PDL * Enrolls mmvf, rope, set-rows and topk kernels for gpt-oss into PDL * Introduce ggml_cuda_kernel_launch, to abstract away cudaLaunchKernelEx, to enable hip/musa compatibility * Enrolls cpy_scalar_contiguous, k_get_rows_float and rms_norm_f32 * Enrolls flash_attn_combine_results * Fix: Drops needless and broken check of CUDA arch for PDL. PDL either works or is without effect. * Enrolls flash-attention kernels to pdl * Fix: inlines ggml_cuda_kernel_launch, and uses perfect forwarding for kernels args. This fixes PDL. * Perf: Enrolls k_bin_bcast variadic template invocation into PDL, via and template alias and template expansion * Enrolls all remaining kernels for qwen3-coder-next into PDL * Remove all PDL LC calls to create a baseline * Added LC according to internal guidance and tested kernel performance. * Enrols missing qwen3-5 kernels passively into PDL. * Kernel optimizations (LC signals) for qwen3.5 * Enrolls ssm-scan kernels into PDL * Adds GGML_CUDA_PDL command line option to toggle PDL. * Fix: Ada and lower compilation by guarding PDL calls correctly * Cleanup: Removes commented out GGML_CUDA_PDL_LC * Cleanup: Removes experimental comments * Adds 90-virtual to build script so that Hopper GPUs can leverage PDL. * Adds stricter checks to enable PDL, adds env-check to disable it, and removes now superfluous compile option to enable PDL. * Fix: Correct PDL en/disablement based on device-side arch check. Host side check is UB. Required moving from macros to inlined functions * Fix: default-disable PDL. Enable by setting GGML_CUDA_ENABLE_PDL=1 * Enable PDL by default for Hopper+ devices * Enrolls softcap_f32 and two flash_attn kernels into PDL. * Improves flash attn PDL barrier placement * Fix: Perf regression on ada; excludes ada and below from PDL launches * Improves some sync barrier placements * Drops superfluous constructor * Adds #endif guard comments * Reverts experimental change to top-k-moe.cu, which moved expensive allocations in front of the PDL barrier. It did not have a meaningful impact. * Exchanges GGML_CUDA_DISABLE_PDL with GGML_CUDA_PDL. IFF GGML_CUDA_PDL=0 PDL is disabled * Revert "Drops superfluous constructor". Adds const to remaining arguments This reverts commit 12b1d250da0089ae02a9bb71bbb3fd6d70f6f2f1. * Cleanup: Removes and fixes some comments and whitespace * Clarifies comment of sync-barrier position * Relocates and refactors PDL launch functions and accessories * Adds error checking to the regular kernel launch path * Drops "auto" in favor of "ggml_cuda_kernel_params" * Adds "const" to ggml_cuda_kernel_launch_params * [Whitespace] Adds final newline to common.cuh to make editorconfig CI job happy
78 lines
3.0 KiB
Plaintext
78 lines
3.0 KiB
Plaintext
#include "mean.cuh"
|
|
#include "reduce_rows.cuh"
|
|
|
|
#ifdef GGML_CUDA_USE_CUB
|
|
#include <cub/cub.cuh>
|
|
using namespace cub;
|
|
#endif // GGML_CUDA_USE_CUB
|
|
|
|
template <typename T> __global__ void divide_by_count(T * result, size_t count) {
|
|
*result /= static_cast<T>(count);
|
|
}
|
|
|
|
void ggml_cuda_op_mean(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
|
|
const ggml_tensor * src0 = dst->src[0];
|
|
const float * src0_d = (const float *) src0->data;
|
|
float * dst_d = (float *) dst->data;
|
|
cudaStream_t stream = ctx.stream();
|
|
|
|
GGML_ASSERT(src0->type == GGML_TYPE_F32);
|
|
GGML_ASSERT(dst->type == GGML_TYPE_F32);
|
|
GGML_ASSERT(ggml_is_contiguous(src0));
|
|
|
|
const int64_t ncols = src0->ne[0];
|
|
const int64_t nrows = ggml_nrows(src0);
|
|
|
|
// Special case for reducing vectors
|
|
#ifdef GGML_CUDA_USE_CUB
|
|
#ifdef USE_CUDA_GRAPH
|
|
cudaStreamCaptureStatus iscapturing;
|
|
CUDA_CHECK(cudaStreamIsCapturing(stream, &iscapturing));
|
|
#endif // USE_CUDA_GRAPH
|
|
if ((nrows == 1) &&
|
|
#ifdef USE_CUDA_GRAPH
|
|
// Determine if CUDA graphs are effectively disabled for this context
|
|
// (no graph instance exists and we're not capturing, OR graphs are explicitly enabled)
|
|
(((ncols > 65536) &&
|
|
(((!ctx.any_cuda_graph_has_instance()) && (iscapturing == cudaStreamCaptureStatusNone)) ||
|
|
ctx.any_cuda_graph_enabled())) ||
|
|
// CUDA graphs are enabled - use lower threshold
|
|
((ncols > 32768) &&
|
|
!(((!ctx.any_cuda_graph_has_instance()) && (iscapturing == cudaStreamCaptureStatusNone)) ||
|
|
ctx.any_cuda_graph_enabled())))) {
|
|
#else
|
|
(ncols > 65536)) {
|
|
#endif // USE_CUDA_GRAPH
|
|
// Single row - use device-wide reduction
|
|
size_t tmp_size = 0;
|
|
ggml_cuda_pool & pool = ctx.pool();
|
|
|
|
DeviceReduce::Sum(nullptr, tmp_size, src0_d, dst_d, ncols, stream);
|
|
|
|
ggml_cuda_pool_alloc<uint8_t> tmp_alloc(pool, tmp_size);
|
|
DeviceReduce::Sum(tmp_alloc.ptr, tmp_size, src0_d, dst_d, ncols, stream);
|
|
|
|
// Divide by ncols
|
|
divide_by_count<float><<<1, 1, 0, stream>>>(dst_d, ncols);
|
|
return;
|
|
}
|
|
#endif // GGML_CUDA_USE_CUB
|
|
|
|
const dim3 block_nums(nrows, 1, 1);
|
|
|
|
const int id = ggml_cuda_get_device();
|
|
const int nsm = ggml_cuda_info().devices[id].nsm;
|
|
|
|
// Heuristic for block size selection to optimize occupancy.
|
|
// See discussion in: https://github.com/ggml-org/llama.cpp/pull/15132
|
|
if ((nrows / nsm) < 2) {
|
|
const dim3 block_dims(512, 1, 1);
|
|
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, 0, stream);
|
|
ggml_cuda_kernel_launch(reduce_rows_f32</*norm=*/true>, launch_params, src0_d, dst_d, ncols);
|
|
} else {
|
|
const dim3 block_dims(ncols < 1024 ? 32 : 128, 1, 1);
|
|
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, 0, stream);
|
|
ggml_cuda_kernel_launch(reduce_rows_f32</*norm=*/true>, launch_params, src0_d, dst_d, ncols);
|
|
}
|
|
}
|