Commit 4b51674 for stable-diffusion.cpp
commit 4b51674b662fe71688a63e06321f09cb3e7915a1
Author: leejet <leejet714@gmail.com>
Date: Sat Nov 14 15:22:31 2026 +0800
perf(hip): accelerate flash attention and INT8 projections (#2351)
diff --git a/docs/build.md b/docs/build.md
index 8cf235b..475800f 100644
--- a/docs/build.md
+++ b/docs/build.md
@@ -165,6 +165,20 @@ hipBLAS, rocBLAS, and their development files. See
[AMD's ROCm installation guide](https://rocm.docs.amd.com/en/latest/install/rocm.html)
for supported GPUs, operating systems, and installation methods.
+When the hipBLASLt development package (version 1.0 or later) is available,
+CMake also enables INT8 GEMM tuning. The first run of an eligible matrix shape
+benchmarks the no-workspace candidates returned by hipBLASLt for the current
+device, including CDNA devices, and caches the result in the process. It
+selects a hipBLASLt candidate that runs successfully and beats hipBLAS, or a
+working candidate if hipBLAS does not support the operation.
+Set `SD_HIP_DISABLE_I8_LT=1` to use the hipBLAS GEMM path for comparison.
+
+With a compiled wave32 WMMA kernel, the HIP flash attention path supports
+head dimensions 40, 52, 64, 72, 80, 120, 128, 160, 256, and 320, including
+grouped-query attention. It requires unmasked attention with at least 512
+query tokens and 512 key tokens. Set `SD_HIP_DISABLE_WMMA_ATTN=1` to use
+ggml's native attention path for comparison.
+
**Linux**
With ROCm installed, use the default CMake generator and host compilers:
diff --git a/src/core/cuda/CMakeLists.txt b/src/core/cuda/CMakeLists.txt
index e2dd3f9..a1290e7 100644
--- a/src/core/cuda/CMakeLists.txt
+++ b/src/core/cuda/CMakeLists.txt
@@ -49,16 +49,27 @@ if(SD_HIPBLAS)
find_package(hipblas REQUIRED)
if(WIN32 OR CMAKE_CXX_COMPILER MATCHES "hipcc")
sd_cuda_sources(hip CXX)
+ set_source_files_properties(hip_attention.cu PROPERTIES LANGUAGE CXX)
target_link_libraries(${SD_CUDA_TARGET} PRIVATE hip::device)
else()
sd_cuda_sources(hip HIP)
+ set_source_files_properties(hip_attention.cu PROPERTIES LANGUAGE HIP)
set_target_properties(${SD_CUDA_TARGET} PROPERTIES HIP_STANDARD 17 HIP_STANDARD_REQUIRED ON)
endif()
+ target_sources(${SD_CUDA_TARGET} PRIVATE hip_attention.cu)
set_source_files_properties(${SD_CUDA_SOURCES} PROPERTIES COMPILE_OPTIONS "-ffast-math")
# W4A8 WMMA uses the wave64 fragment layout on both GFX11 and GFX12.
set_property(SOURCE "${CMAKE_CURRENT_BINARY_DIR}/w4_hip.cu" APPEND PROPERTY COMPILE_OPTIONS "-mwavefrontsize64")
target_link_libraries(${SD_CUDA_TARGET} PRIVATE hip::host roc::hipblas)
target_link_libraries(${SD_LIB} PRIVATE hip::host roc::hipblas)
+ find_package(hipblaslt CONFIG QUIET)
+ if(hipblaslt_FOUND AND TARGET roc::hipblaslt AND hipblaslt_VERSION VERSION_GREATER_EQUAL 1.0)
+ target_sources(${SD_CUDA_TARGET} PRIVATE hip_i8_gemm.cpp)
+ target_include_directories(${SD_CUDA_TARGET} PRIVATE "${PROJECT_SOURCE_DIR}/include")
+ target_compile_definitions(${SD_CUDA_TARGET} PRIVATE SD_HIP_I8_LT)
+ target_link_libraries(${SD_CUDA_TARGET} PRIVATE roc::hipblaslt)
+ target_link_libraries(${SD_LIB} PRIVATE roc::hipblaslt)
+ endif()
endif()
if(SD_MUSA)
diff --git a/src/core/cuda/convrot_kernel.h b/src/core/cuda/convrot_kernel.h
index fc0d147..fa9e4da 100644
--- a/src/core/cuda/convrot_kernel.h
+++ b/src/core/cuda/convrot_kernel.h
@@ -125,8 +125,8 @@ namespace SD_CUDA_NAMESPACE {
if constexpr (dequantize) {
gate_value = (float)((const int32_t*)src_row)[i] * activation_scales[row];
up_value = (float)((const int32_t*)up)[row * up_stride + i] * activation_scales[row];
- gate_value *= weight_scales[i];
- up_value *= weight_scales[k + i];
+ gate_value *= weight_scales[(situ_gate ? k : 0) + i];
+ up_value *= weight_scales[(situ_gate ? 0 : k) + i];
} else {
gate_value = src_row[i];
up_value = up[row * up_stride + i];
diff --git a/src/core/cuda/hip_attention.cu b/src/core/cuda/hip_attention.cu
new file mode 100644
index 0000000..adff9fb
--- /dev/null
+++ b/src/core/cuda/hip_attention.cu
@@ -0,0 +1,390 @@
+#define SD_CUDA_VENDOR_HIP 1
+#include "runtime.h"
+
+#include <climits>
+#include <cmath>
+#include <cstdlib>
+#include <cstring>
+
+namespace sd_hip {
+ static constexpr int WAVE_SIZE = 32;
+ static constexpr int WMMA_TILE = 16;
+
+ template <int HeadDim>
+ struct AttentionTile {
+ static constexpr int channels = (HeadDim + WMMA_TILE - 1) / WMMA_TILE * WMMA_TILE;
+ static constexpr int keys = HeadDim > 128 ? 32 : 64;
+ static constexpr int query_waves = HeadDim > 128 ? 4 : 8;
+ static constexpr int queries = query_waves * WMMA_TILE;
+ };
+
+#if __has_builtin(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12)
+#define SD_HIP_ATTN_WMMA_128B
+#elif __has_builtin(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32)
+#define SD_HIP_ATTN_WMMA_256B
+#endif
+
+#if defined(SD_HIP_ATTN_WMMA_128B) || defined(SD_HIP_ATTN_WMMA_256B)
+ using Half8 = _Float16 __attribute__((ext_vector_type(8)));
+ using Accumulator = float __attribute__((ext_vector_type(8)));
+#ifdef SD_HIP_ATTN_WMMA_256B
+ using HalfFragment = _Float16 __attribute__((ext_vector_type(16)));
+ static constexpr bool replicate_inputs = true;
+#else
+ using HalfFragment = Half8;
+ static constexpr bool replicate_inputs = false;
+#endif
+
+ static __device__ __forceinline__ HalfFragment load_fragment(const _Float16* ptr) {
+ const auto low = *reinterpret_cast<const Half8*>(ptr);
+#ifdef SD_HIP_ATTN_WMMA_256B
+ // LDS rows are only 16-byte aligned, including when the fragment spans 32 bytes.
+ const auto high = *reinterpret_cast<const Half8*>(ptr + 8);
+ return __builtin_shufflevector(low, high, 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
+#else
+ return low;
+#endif
+ }
+
+ static __device__ __forceinline__ Accumulator attention_wmma(HalfFragment a, HalfFragment b, Accumulator acc) {
+#ifdef SD_HIP_ATTN_WMMA_256B
+ return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, acc);
+#else
+ return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(a, b, acc);
+#endif
+ }
+#endif
+
+ struct AttentionInput {
+ const void* data;
+ size_t token_stride;
+ size_t head_stride;
+ size_t batch_stride;
+
+ template <typename T>
+ __device__ const T* head(int index, int batch) const {
+ return reinterpret_cast<const T*>(static_cast<const char*>(data) + index * head_stride + batch * batch_stride);
+ }
+ };
+
+ template <int HeadDim>
+ static __global__ void pack_attention_values(AttentionInput value, _Float16* packed, int keys) {
+ constexpr int channels = AttentionTile<HeadDim>::channels;
+ constexpr int key_tile = AttentionTile<HeadDim>::keys;
+ constexpr int channel_tiles = (HeadDim + 31) / 32;
+ // Each block transposes a 32x32 tile. The extra column avoids LDS bank conflicts.
+ __shared__ _Float16 tile[32][33];
+ const int head = blockIdx.y;
+ const int batch = blockIdx.z;
+ const int token = (blockIdx.x / channel_tiles) * 32;
+ const int channel = (blockIdx.x % channel_tiles) * 32;
+ const int x = threadIdx.x;
+ const int y = threadIdx.y;
+ const auto* input = value.head<_Float16>(head, batch);
+ const size_t stride = value.token_stride / sizeof(_Float16);
+ const size_t head_size = ((size_t(keys) + key_tile - 1) / key_tile) * key_tile * channels;
+ auto* output = packed + (size_t(batch) * gridDim.y + head) * head_size +
+ size_t(token / key_tile) * key_tile * channels + channel * key_tile + token % key_tile;
+#pragma unroll
+ for (int row = y; row < 32; row += 8) {
+ tile[row][x] = token + row < keys && channel + x < HeadDim ? input[(token + row) * stride + channel + x] : _Float16(0);
+ }
+ __syncthreads();
+#pragma unroll
+ for (int row = y; row < 32; row += 8) {
+ if (channel + row < channels)
+ output[row * key_tile + x] = tile[x][row];
+ }
+ }
+
+ template <int HeadDim>
+ static __global__ __launch_bounds__(AttentionTile<HeadDim>::query_waves* WAVE_SIZE) void wmma_attention(
+ AttentionInput query,
+ AttentionInput key,
+ const _Float16* packed_value,
+ float* dst,
+ int queries,
+ int keys,
+ int kv_heads,
+ int queries_per_kv,
+ float scale) {
+#if defined(SD_HIP_ATTN_WMMA_128B) || defined(SD_HIP_ATTN_WMMA_256B)
+ constexpr int channels = AttentionTile<HeadDim>::channels;
+ constexpr int key_tile = AttentionTile<HeadDim>::keys;
+ constexpr int query_tile = AttentionTile<HeadDim>::queries;
+ const int lane = threadIdx.x % WMMA_TILE;
+ const int half = (threadIdx.x % WAVE_SIZE) / WMMA_TILE;
+ // The 256-bit operand layout duplicates both fragments across the two half-waves
+ // and distributes accumulator rows as even and odd indices.
+ const int input_offset = replicate_inputs ? 0 : half * 8;
+ const int acc_offset = replicate_inputs ? half : half * 8;
+ constexpr int acc_stride = replicate_inputs ? 2 : 1;
+ const int wave = threadIdx.x / WAVE_SIZE;
+ const int head = blockIdx.y;
+ const int kv_head = queries_per_kv == 1 ? head : head / queries_per_kv;
+ const int batch = blockIdx.z;
+ const int q_index = blockIdx.x * query_tile + wave * WMMA_TILE + lane;
+ const float* q = query.head<float>(head, batch);
+ const _Float16* k = key.head<_Float16>(kv_head, batch);
+ const size_t q_stride = query.token_stride / sizeof(float);
+ const size_t k_stride = key.token_stride / sizeof(_Float16);
+ const size_t value_head_size = ((size_t(keys) + key_tile - 1) / key_tile) * key_tile * channels;
+ const _Float16* v = packed_value + (size_t(batch) * kv_heads + kv_head) * value_head_size;
+
+ __shared__ __align__(16) _Float16 k_tile[key_tile][channels + 8];
+ HalfFragment q_fragments[channels / WMMA_TILE] = {};
+ Accumulator output[channels / WMMA_TILE] = {};
+ if (q_index < queries) {
+#pragma unroll
+ for (int d = 0; d < channels / WMMA_TILE; ++d) {
+#pragma unroll
+ for (int i = 0; i < (replicate_inputs ? 16 : 8); ++i) {
+ const int channel = d * WMMA_TILE + input_offset + i;
+ if (channel < HeadDim)
+ q_fragments[d][i] = q[q_index * q_stride + channel] * scale;
+ }
+ }
+ }
+
+ float maximum = -INFINITY;
+ float sum = 0.f;
+ for (int base = 0; base < keys; base += key_tile) {
+ const int tile_keys = min(keys - base, key_tile);
+ for (int index = threadIdx.x; index < key_tile * channels / 8; index += blockDim.x) {
+ const int token = index / (channels / 8);
+ const int channel = (index % (channels / 8)) * 8;
+ Half8 k_fragment = {};
+ if (token < tile_keys && channel < HeadDim) {
+ const auto* row = k + (base + token) * k_stride;
+ if constexpr (HeadDim % 8 == 0) {
+ k_fragment = *reinterpret_cast<const Half8*>(row + channel);
+ } else {
+#pragma unroll
+ for (int i = 0; i < 8; ++i)
+ if (channel + i < HeadDim)
+ k_fragment[i] = row[channel + i];
+ }
+ }
+ *reinterpret_cast<Half8*>(&k_tile[token][channel]) = k_fragment;
+ }
+ __syncthreads();
+
+ // K * Q^T leaves each query in a pair of lanes.
+ Accumulator scores[key_tile / WMMA_TILE] = {};
+#pragma unroll
+ for (int d = 0; d < channels / WMMA_TILE; ++d) {
+#pragma unroll
+ for (int t = 0; t < key_tile / WMMA_TILE; ++t) {
+ const auto k_fragment = load_fragment(&k_tile[t * WMMA_TILE + lane][d * WMMA_TILE + input_offset]);
+ scores[t] = attention_wmma(k_fragment, q_fragments[d], scores[t]);
+ }
+ }
+
+ float tile_max = -INFINITY;
+#pragma unroll
+ for (int t = 0; t < key_tile / WMMA_TILE; ++t) {
+#pragma unroll
+ for (int i = 0; i < 8; ++i) {
+ scores[t][i] = t * WMMA_TILE + acc_offset + acc_stride * i < tile_keys ? scores[t][i] : -INFINITY;
+ tile_max = fmaxf(tile_max, scores[t][i]);
+ }
+ }
+ const float next_max = fmaxf(maximum, fmaxf(tile_max, cuda_shuffle_xor(tile_max, 16)));
+ const float rescale = __expf(maximum - next_max);
+ maximum = next_max;
+ float tile_sum = 0.f;
+ HalfFragment probabilities[key_tile / WMMA_TILE];
+#pragma unroll
+ for (int t = 0; t < key_tile / WMMA_TILE; ++t) {
+#pragma unroll
+ for (int i = 0; i < 8; ++i) {
+ const float p = __expf(scores[t][i] - maximum);
+ if constexpr (replicate_inputs) {
+ // Reassemble the 16 probabilities in key order for both
+ // half-waves before using them as the B operand of V^T * P.
+ const float peer = cuda_shuffle_xor(p, WMMA_TILE);
+ probabilities[t][2 * i] = half == 0 ? p : peer;
+ probabilities[t][2 * i + 1] = half == 0 ? peer : p;
+ } else {
+ probabilities[t][i] = p;
+ }
+ tile_sum += p;
+ }
+ }
+ sum = fmaf(sum, rescale, tile_sum + cuda_shuffle_xor(tile_sum, 16));
+#pragma unroll
+ for (int d = 0; d < channels / WMMA_TILE; ++d) {
+ output[d] *= rescale;
+#pragma unroll
+ for (int t = 0; t < key_tile / WMMA_TILE; ++t) {
+ const size_t offset = size_t(base) * channels + (d * WMMA_TILE + lane) * key_tile + t * WMMA_TILE + input_offset;
+ const auto v_fragment = load_fragment(v + offset);
+ output[d] = attention_wmma(v_fragment, probabilities[t], output[d]);
+ }
+ }
+ __syncthreads();
+ }
+
+ if (q_index < queries) {
+ const float inverse = 1.f / sum;
+ float* out = dst + ((size_t(batch) * queries + q_index) * gridDim.y + head) * HeadDim;
+#pragma unroll
+ for (int d = 0; d < channels / WMMA_TILE; ++d) {
+ output[d] *= inverse;
+ if constexpr (replicate_inputs || HeadDim % 8 != 0) {
+#pragma unroll
+ for (int i = 0; i < 8; ++i) {
+ const int channel = d * WMMA_TILE + acc_offset + acc_stride * i;
+ if (channel < HeadDim)
+ out[channel] = output[d][i];
+ }
+ } else {
+ const int channel = d * WMMA_TILE + acc_offset;
+ if (channel < HeadDim)
+ *reinterpret_cast<Accumulator*>(out + channel) = output[d];
+ }
+ }
+ }
+#elif defined(__HIP_DEVICE_COMPILE__)
+ __builtin_trap();
+#endif
+ }
+
+ static bool attention_aligned(const ggml_tensor* tensor, size_t alignment) {
+ size_t offset = 0;
+ while (tensor->view_src) {
+ offset += tensor->view_offs;
+ tensor = tensor->view_src;
+ }
+ return (reinterpret_cast<uintptr_t>(tensor->data) + offset) % alignment == 0;
+ }
+
+ static bool attention_layout(const ggml_tensor* tensor, ggml_type type, size_t alignment) {
+ return tensor->type == type && tensor->nb[0] == ggml_type_size(type) &&
+ tensor->nb[1] % alignment == 0 && tensor->nb[2] % alignment == 0 && tensor->nb[3] % alignment == 0 &&
+ attention_aligned(tensor, alignment);
+ }
+
+ template <int HeadDim>
+ static ggml_status launch_attention(const ggml_ext_execution_context* execution, ggml_tensor* const* nodes, CUDAStream stream) {
+ using Tile = AttentionTile<HeadDim>;
+ auto* dst = nodes[0];
+ auto* q = dst->src[0];
+ auto* k = dst->src[1];
+ auto* v = dst->src[2];
+ auto input = [](const ggml_tensor* t) { return AttentionInput{t->data, t->nb[1], t->nb[2], t->nb[3]}; };
+ float scale;
+ std::memcpy(&scale, dst->op_params, sizeof(scale));
+ auto* packed = static_cast<_Float16*>(ggml_backend_buffer_get_base(execution->workspace));
+ const size_t value_tiles = ((v->ne[1] + Tile::keys - 1) / Tile::keys) * (Tile::keys / 32) * ((HeadDim + 31) / 32);
+ const dim3 pack_grid(value_tiles, v->ne[2], v->ne[3]);
+ pack_attention_values<HeadDim><<<pack_grid, dim3(32, 8), 0, stream>>>(input(v), packed, v->ne[1]);
+ if (cuda_last_error() != cuda_success)
+ return GGML_STATUS_FAILED;
+
+ const dim3 grid((q->ne[1] + Tile::queries - 1) / Tile::queries, q->ne[2], q->ne[3]);
+ wmma_attention<HeadDim><<<grid, Tile::query_waves * WAVE_SIZE, 0, stream>>>(
+ input(q), input(k), packed, static_cast<float*>(dst->data), q->ne[1], k->ne[1], k->ne[2], q->ne[2] / k->ne[2], scale);
+ return cuda_last_error() == cuda_success ? GGML_STATUS_SUCCESS : GGML_STATUS_FAILED;
+ }
+
+ struct AttentionKernel {
+ int head_dim;
+ int channels;
+ int key_tile;
+ int query_tile;
+ ggml_status (*launch)(const ggml_ext_execution_context*, ggml_tensor* const*, CUDAStream);
+ };
+
+ template <int HeadDim>
+ static constexpr AttentionKernel attention_kernel() {
+ using Tile = AttentionTile<HeadDim>;
+ return {HeadDim, Tile::channels, Tile::keys, Tile::queries, launch_attention<HeadDim>};
+ }
+
+ static const AttentionKernel* select_attention_kernel(int64_t head_dim) {
+ static constexpr AttentionKernel kernels[] = {
+ attention_kernel<40>(),
+ attention_kernel<52>(),
+ attention_kernel<64>(),
+ attention_kernel<72>(),
+ attention_kernel<80>(),
+ attention_kernel<120>(),
+ attention_kernel<128>(),
+ attention_kernel<160>(),
+ attention_kernel<256>(),
+ attention_kernel<320>(),
+ };
+ for (const auto& kernel : kernels)
+ if (head_dim == kernel.head_dim)
+ return &kernel;
+ return nullptr;
+ }
+
+ static bool wmma_attention_supports(ggml_backend_dev_t, const ggml_tensor* dst, void*) {
+ const auto* q = dst->src[0];
+ const auto* k = dst->src[1];
+ const auto* v = dst->src[2];
+ const auto* kernel = select_attention_kernel(q->ne[0]);
+ float params[3];
+ std::memcpy(params, dst->op_params, sizeof(params));
+ return kernel && !dst->src[3] && !dst->src[4] && std::isfinite(params[0]) && params[1] == 0.f && params[2] == 0.f &&
+ q->ne[0] == k->ne[0] && ggml_are_same_shape(k, v) &&
+ k->ne[2] > 0 && q->ne[2] % k->ne[2] == 0 && q->ne[3] == k->ne[3] && q->ne[2] <= 65535 && q->ne[3] <= 65535 &&
+ q->ne[1] >= 512 && k->ne[1] >= 512 && q->ne[1] <= INT_MAX - kernel->query_tile &&
+ k->ne[1] <= INT_MAX - kernel->key_tile &&
+ attention_layout(q, GGML_TYPE_F32, sizeof(float)) &&
+ attention_layout(k, GGML_TYPE_F16, q->ne[0] % 8 == 0 ? 16 : sizeof(_Float16)) &&
+ attention_layout(v, GGML_TYPE_F16, sizeof(_Float16)) &&
+ dst->type == GGML_TYPE_F32 && ggml_is_contiguous(dst) && !dst->view_src && attention_aligned(dst, 32);
+ }
+
+ static ggml_status wmma_attention_prepare(ggml_backend_t, ggml_tensor* const* nodes, int, void**, size_t* workspace, void*) {
+ const auto* v = nodes[0]->src[2];
+ const auto* kernel = select_attention_kernel(v->ne[0]);
+ if (!kernel)
+ return GGML_STATUS_FAILED;
+ const size_t tile = kernel->key_tile;
+ *workspace = ((size_t(v->ne[1]) + tile - 1) / tile) * tile * kernel->channels *
+ v->ne[2] * v->ne[3] * sizeof(_Float16);
+ return GGML_STATUS_SUCCESS;
+ }
+
+ static void wmma_attention_destroy(void*, void*) {}
+
+ static ggml_status wmma_attention_execute(const ggml_ext_execution_context* execution, ggml_tensor* const* nodes, int, void*, void*) {
+ ggml_ext_gpu_context context;
+ if (!execution->workspace || !cuda_context(execution->backend, &context)) {
+ return GGML_STATUS_FAILED;
+ }
+ const auto* kernel = select_attention_kernel(nodes[0]->src[0]->ne[0]);
+ return kernel ? kernel->launch(execution, nodes, static_cast<CUDAStream>(context.stream)) : GGML_STATUS_FAILED;
+ }
+
+ void register_wmma_attention(ggml_backend_dev_t device) {
+ static const bool disabled = std::getenv("SD_HIP_DISABLE_WMMA_ATTN") != nullptr;
+ const auto info = cuda_device_info(device);
+ int previous;
+ if (disabled || info.device < 0 || hipGetDevice(&previous) != hipSuccess || hipSetDevice(info.device) != hipSuccess) {
+ return;
+ }
+ hipFuncAttributes attributes{};
+ const auto status = hipFuncGetAttributes(&attributes, reinterpret_cast<const void*>(wmma_attention<128>));
+ const auto restored = hipSetDevice(previous);
+ // The device image uses a trap stub with no LDS when it lacks compatible wave32 WMMA.
+ if (status != hipSuccess || restored != hipSuccess || attributes.sharedSizeBytes == 0) {
+ return;
+ }
+ ggml_ext_kernel_desc desc = {};
+ desc.abi_version = GGML_EXT_ABI_VERSION;
+ desc.name = "sd.wmma_attention.hip.v1";
+ desc.op = GGML_OP_FLASH_ATTN_EXT;
+ desc.priority = 200;
+ desc.supports = wmma_attention_supports;
+ desc.preferred = wmma_attention_supports;
+ desc.prepare = wmma_attention_prepare;
+ desc.execute = wmma_attention_execute;
+ desc.destroy = wmma_attention_destroy;
+ GGML_ASSERT(ggml_ext_register_kernel_ex(device, &desc, GGML_EXT_KERNEL_CAPTURE_SAFE | GGML_EXT_KERNEL_REPLAY_REUSABLE));
+ }
+}
diff --git a/src/core/cuda/hip_i8_gemm.cpp b/src/core/cuda/hip_i8_gemm.cpp
new file mode 100644
index 0000000..61b5fbb
--- /dev/null
+++ b/src/core/cuda/hip_i8_gemm.cpp
@@ -0,0 +1,220 @@
+#include "hip_i8_gemm.h"
+
+#include <hipblaslt/hipblaslt.h>
+#include <algorithm>
+#include <array>
+#include <cstdlib>
+#include <limits>
+#include <list>
+#include <map>
+#include <mutex>
+#include <tuple>
+
+#include "core/util.h"
+
+namespace sd_hip {
+ using GemmShape = std::tuple<int, int, int>;
+
+ struct HipLtPlan {
+ GemmShape shape;
+ hipblasLtMatmulDesc_t op = nullptr;
+ hipblasLtMatrixLayout_t a = nullptr;
+ hipblasLtMatrixLayout_t b = nullptr;
+ hipblasLtMatrixLayout_t c = nullptr;
+ hipblasLtMatmulAlgo_t algo{};
+ bool use_lt = false;
+
+ explicit HipLtPlan(GemmShape shape)
+ : shape(shape) {}
+
+ ~HipLtPlan() {
+ if (c)
+ hipblasLtMatrixLayoutDestroy(c);
+ if (b)
+ hipblasLtMatrixLayoutDestroy(b);
+ if (a)
+ hipblasLtMatrixLayoutDestroy(a);
+ if (op)
+ hipblasLtMatmulDescDestroy(op);
+ }
+
+ bool init(const HipI8Gemm& gemm) {
+ const hipblasOperation_t transpose = HIPBLAS_OP_T;
+ return hipblasLtMatmulDescCreate(&op, HIPBLAS_COMPUTE_32I, HIP_R_32I) == HIPBLAS_STATUS_SUCCESS &&
+ hipblasLtMatmulDescSetAttribute(op, HIPBLASLT_MATMUL_DESC_TRANSA, &transpose, sizeof(transpose)) == HIPBLAS_STATUS_SUCCESS &&
+ hipblasLtMatrixLayoutCreate(&a, HIP_R_8I, gemm.k, gemm.n, gemm.k) == HIPBLAS_STATUS_SUCCESS &&
+ hipblasLtMatrixLayoutCreate(&b, HIP_R_8I, gemm.k, gemm.rows, gemm.k) == HIPBLAS_STATUS_SUCCESS &&
+ hipblasLtMatrixLayoutCreate(&c, HIP_R_32I, gemm.n, gemm.rows, gemm.n) == HIPBLAS_STATUS_SUCCESS;
+ }
+
+ hipblasStatus_t run(hipblasLtHandle_t handle, const HipI8Gemm& gemm, const hipblasLtMatmulAlgo_t& choice) const {
+ const int32_t alpha = 1, beta = 0;
+ return hipblasLtMatmul(handle, op, &alpha, gemm.weight, a, gemm.input, b,
+ &beta, gemm.output, c, gemm.output, c, &choice, nullptr, 0, gemm.stream);
+ }
+ };
+
+ struct HipLtDevice {
+ int ordinal;
+ std::mutex mutex;
+ std::map<hipStream_t, hipblasLtHandle_t> handles;
+ std::list<HipLtPlan> plans;
+
+ explicit HipLtDevice(int ordinal)
+ : ordinal(ordinal) {}
+
+ ~HipLtDevice() {
+ if (handles.empty())
+ return;
+ int previous;
+ if (hipGetDevice(&previous) == hipSuccess && hipSetDevice(ordinal) == hipSuccess) {
+ for (const auto& entry : handles)
+ (void)hipblasLtDestroy(entry.second);
+ (void)hipSetDevice(previous);
+ }
+ }
+ };
+
+ struct HipGemmTimer {
+ hipEvent_t start = nullptr;
+ hipEvent_t end = nullptr;
+
+ ~HipGemmTimer() {
+ if (end)
+ (void)hipEventDestroy(end);
+ if (start)
+ (void)hipEventDestroy(start);
+ }
+
+ template <typename Launch>
+ float measure(hipStream_t stream, Launch launch) const {
+ constexpr int trials = 3;
+ if (hipEventRecord(start, stream) != hipSuccess)
+ return 0.f;
+ for (int i = 0; i < trials; ++i) {
+ if (launch() != HIPBLAS_STATUS_SUCCESS)
+ return 0.f;
+ }
+ float ms;
+ return hipEventRecord(end, stream) == hipSuccess && hipEventSynchronize(end) == hipSuccess &&
+ hipEventElapsedTime(&ms, start, end) == hipSuccess
+ ? ms / trials
+ : 0.f;
+ }
+ };
+
+ static void tune(int device, hipblasLtHandle_t handle, const HipI8Gemm& gemm, HipLtPlan& plan) {
+ if (!plan.init(gemm))
+ return;
+
+ hipblasLtMatmulPreference_t preference = nullptr;
+ if (hipblasLtMatmulPreferenceCreate(&preference) != HIPBLAS_STATUS_SUCCESS)
+ return;
+ // A zero-workspace kernel can use the existing ggml accumulator buffer.
+ constexpr uint64_t workspace = 0;
+ std::array<hipblasLtMatmulHeuristicResult_t, 8> choices{};
+ int count = 0;
+ if (hipblasLtMatmulPreferenceSetAttribute(preference, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &workspace, sizeof(workspace)) == HIPBLAS_STATUS_SUCCESS) {
+ const auto status = hipblasLtMatmulAlgoGetHeuristic(handle, plan.op, plan.a, plan.b, plan.c, plan.c,
+ preference, int(choices.size()), choices.data(), &count);
+ if (status != HIPBLAS_STATUS_SUCCESS)
+ count = 0;
+ }
+ hipblasLtMatmulPreferenceDestroy(preference);
+ if (count == 0)
+ return;
+
+ HipGemmTimer timer;
+ if (hipEventCreate(&timer.start) != hipSuccess || hipEventCreate(&timer.end) != hipSuccess)
+ return;
+ const auto blas_status = gemm.run_hipblas();
+ const bool blas_unsupported = blas_status == HIPBLAS_STATUS_NOT_SUPPORTED || blas_status == HIPBLAS_STATUS_ARCH_MISMATCH;
+ float baseline = 0.f;
+ if (blas_status == HIPBLAS_STATUS_SUCCESS) {
+ baseline = timer.measure(gemm.stream, [&] { return gemm.run_hipblas(); });
+ if (baseline <= 0.f)
+ return;
+ } else if (!blas_unsupported) {
+ return;
+ }
+
+ float best = blas_unsupported ? std::numeric_limits<float>::infinity() : baseline * 0.95f;
+ for (int i = 0; i < count; ++i) {
+ const auto& candidate = choices[i];
+ if (candidate.state != HIPBLAS_STATUS_SUCCESS || candidate.workspaceSize != 0 ||
+ plan.run(handle, gemm, candidate.algo) != HIPBLAS_STATUS_SUCCESS)
+ continue;
+ const float ms = timer.measure(gemm.stream, [&] { return plan.run(handle, gemm, candidate.algo); });
+ if (ms > 0.f && ms < best) {
+ best = ms;
+ plan.algo = candidate.algo;
+ plan.use_lt = true;
+ }
+ }
+ if (plan.use_lt) {
+ if (blas_unsupported) {
+ LOG_VERBOSE("hipBLASLt INT8 GEMM %d x %d x %d on HIP device %d: %.3f ms (hipBLAS unsupported)",
+ gemm.n, gemm.rows, gemm.k, device, best);
+ } else {
+ LOG_VERBOSE("hipBLASLt INT8 GEMM %d x %d x %d on HIP device %d: %.3f -> %.3f ms",
+ gemm.n, gemm.rows, gemm.k, device, baseline, best);
+ }
+ }
+ }
+
+ static HipLtDevice& lt_device(int ordinal) {
+ static std::mutex mutex;
+ static std::map<int, HipLtDevice> devices;
+ std::lock_guard<std::mutex> lock(mutex);
+ return devices.try_emplace(ordinal, ordinal).first->second;
+ }
+
+ bool hip_i8_gemm_lt(int device, const HipI8Gemm& gemm) {
+ static const bool disabled = std::getenv("SD_HIP_DISABLE_I8_LT") != nullptr;
+ if (disabled || gemm.n < 1024 || gemm.rows < 1024 || gemm.k < 1024 || gemm.k > INT32_MAX / 16384 ||
+ (gemm.n | gemm.k) % 16 != 0 ||
+ (reinterpret_cast<uintptr_t>(gemm.weight) | reinterpret_cast<uintptr_t>(gemm.input) |
+ reinterpret_cast<uintptr_t>(gemm.output)) % 16 != 0) {
+ return false;
+ }
+ auto& state = lt_device(device);
+ std::lock_guard<std::mutex> lock(state.mutex);
+ const GemmShape shape(gemm.n, gemm.rows, gemm.k);
+ auto plan = std::find_if(state.plans.begin(), state.plans.end(), [&](const auto& entry) { return entry.shape == shape; });
+ const bool needs_tuning = plan == state.plans.end();
+ if (!needs_tuning) {
+ state.plans.splice(state.plans.begin(), state.plans, plan);
+ if (!plan->use_lt)
+ return false;
+ }
+
+ const auto stream_handle = state.handles.find(gemm.stream);
+ hipblasLtHandle_t handle = stream_handle == state.handles.end() ? nullptr : stream_handle->second;
+ if (needs_tuning || !handle) {
+ // Handle creation and tuning must run outside graph capture.
+ hipStreamCaptureStatus capture;
+ if (hipStreamIsCapturing(gemm.stream, &capture) != hipSuccess || capture != hipStreamCaptureStatusNone)
+ return false;
+ }
+ if (!handle) {
+ // A host lock does not wait for work submitted on another stream.
+ if (hipblasLtCreate(&handle) != HIPBLAS_STATUS_SUCCESS)
+ return false;
+ state.handles.emplace(gemm.stream, handle);
+ }
+
+ if (needs_tuning) {
+ state.plans.emplace_front(shape);
+ plan = state.plans.begin();
+ tune(device, handle, gemm, *plan);
+ if (state.plans.size() > 64)
+ state.plans.pop_back();
+ if (!plan->use_lt)
+ return false;
+ }
+ if (plan->run(handle, gemm, plan->algo) == HIPBLAS_STATUS_SUCCESS)
+ return true;
+ plan->use_lt = false;
+ return false;
+ }
+}
diff --git a/src/core/cuda/hip_i8_gemm.h b/src/core/cuda/hip_i8_gemm.h
new file mode 100644
index 0000000..1c24b32
--- /dev/null
+++ b/src/core/cuda/hip_i8_gemm.h
@@ -0,0 +1,35 @@
+#ifndef __SD_CORE_CUDA_HIP_I8_GEMM_H__
+#define __SD_CORE_CUDA_HIP_I8_GEMM_H__
+
+#include <hip/hip_runtime_api.h>
+#include <hipblas/hipblas.h>
+#include <cstdint>
+
+namespace sd_hip {
+ struct HipI8Gemm {
+ int n, rows, k;
+ const void* weight;
+ const int8_t* input;
+ int32_t* output;
+ hipblasHandle_t blas;
+ hipStream_t stream;
+
+ hipblasStatus_t run_hipblas() const {
+#if HIP_VERSION >= 60500000
+ constexpr auto compute_type = HIPBLAS_COMPUTE_32I;
+#else
+ constexpr auto compute_type = HIPBLAS_R_32I;
+#endif
+ const int32_t alpha = 1, beta = 0;
+ return hipblasGemmEx(blas, HIPBLAS_OP_T, HIPBLAS_OP_N, n, rows, k, &alpha,
+ weight, HIPBLAS_R_8I, k, input, HIPBLAS_R_8I, k,
+ &beta, output, HIPBLAS_R_32I, n, compute_type, HIPBLAS_GEMM_DEFAULT);
+ }
+ };
+
+#ifdef SD_HIP_I8_LT
+ bool hip_i8_gemm_lt(int device, const HipI8Gemm& gemm);
+#endif
+}
+
+#endif
diff --git a/src/core/cuda/int8.cu b/src/core/cuda/int8.cu
index 13203b5..61991b8 100644
--- a/src/core/cuda/int8.cu
+++ b/src/core/cuda/int8.cu
@@ -9,7 +9,7 @@ struct SDI8CudaGemmPlan;
#include "runtime.h"
#if defined(SD_CUDA_VENDOR_HIP)
-#include <hipblas/hipblas.h>
+#include "hip_i8_gemm.h"
#elif !defined(SD_CUDA_VENDOR_MUSA)
#include <cublas_v2.h>
#endif
@@ -362,6 +362,39 @@ namespace SD_CUDA_NAMESPACE {
dst[index] = value;
}
+#if defined(SD_CUDA_VENDOR_HIP)
+ static __device__ __forceinline__ float dequantize_gated_residual(int32_t accum, float row_scale, float weight_scale, float gate, float residual) {
+ const float value = __fmul_rn(__fmul_rn(float(accum), row_scale), weight_scale);
+ return __fadd_rn(__fmul_rn(value, gate), residual);
+ }
+
+ static __global__ void dequantize_gated_residual_vec4(
+ const int4* accum,
+ const float* scales,
+ const float4* weight_scales,
+ const float4* gates,
+ const float4* residual,
+ float4* dst,
+ int64_t n4,
+ int64_t rows) {
+ const int64_t col = int64_t(blockIdx.x) * blockDim.x + threadIdx.x;
+ const int64_t row = blockIdx.y + int64_t(blockIdx.z) * 65535;
+ if (col >= n4 || row >= rows)
+ return;
+
+ const int64_t index = row * n4 + col;
+ const int4 a = accum[index];
+ const float scale = scales[row];
+ const float4 w = weight_scales[col];
+ const float4 g = gates[col];
+ const float4 r = residual[index];
+ dst[index] = make_float4(dequantize_gated_residual(a.x, scale, w.x, g.x, r.x),
+ dequantize_gated_residual(a.y, scale, w.y, g.y, r.y),
+ dequantize_gated_residual(a.z, scale, w.z, g.z, r.z),
+ dequantize_gated_residual(a.w, scale, w.w, g.w, r.w));
+ }
+#endif
+
static bool aligned(const ggml_tensor* tensor, size_t alignment = 4) {
if (!tensor)
return true;
@@ -439,6 +472,7 @@ namespace SD_CUDA_NAMESPACE {
enum Kind { MAT,
QUANTIZE,
GATED,
+ GATED_RESIDUAL,
SWIGLU_QUANTIZE,
CONCAT };
static bool fusion_disabled() {
@@ -452,6 +486,29 @@ namespace SD_CUDA_NAMESPACE {
(op.input->type == GGML_TYPE_F32 || (aligned(op.input, 16) && aligned(op.input_scale, 16)));
}
#endif
+ static const ggml_tensor* residual_input(const ggml_tensor* mul, const ggml_tensor* add) {
+ return add->src[0] == mul ? add->src[1] : add->src[1] == mul ? add->src[0]
+ : nullptr;
+ }
+
+#if defined(SD_CUDA_VENDOR_HIP)
+ static bool match_gated_residual(const ggml_tensor* mat, const ggml_tensor* mul, const ggml_tensor* add, const SDI8Operation& op) {
+ if (mul->op != GGML_OP_MUL || mul->src[0] != mat || add->op != GGML_OP_ADD || !aligned(op.weight_scale, 16))
+ return false;
+ const auto* gate = mul->src[1];
+ const auto* residual = residual_input(mul, add);
+ if (!gate || gate == mat || gate->ne[0] != mat->ne[0] || ggml_nelements(gate) != mat->ne[0] ||
+ !residual || residual == mat || residual == mul)
+ return false;
+ for (const auto* tensor : {mat, mul, add, gate, residual}) {
+ if (tensor->type != GGML_TYPE_F32 || !ggml_is_contiguous(tensor) || !aligned(tensor, 16) ||
+ (tensor != gate && !ggml_are_same_shape(mat, tensor)))
+ return false;
+ }
+ return true;
+ }
+#endif
+
static int match(ggml_backend_dev_t device, const ggml_cgraph* graph, int first, void*) {
auto* g = const_cast<ggml_cgraph*>(graph);
if (fusion_disabled() || first + 1 >= ggml_graph_n_nodes(g))
@@ -461,6 +518,10 @@ namespace SD_CUDA_NAMESPACE {
if (!supports(device, mat, nullptr) || !sd_ggml_i8_describe(mat, &op) || !op.weight_scale || op.bias)
return 0;
const auto* next = ggml_graph_node(g, first + 1);
+#if defined(SD_CUDA_VENDOR_HIP)
+ if (first + 2 < ggml_graph_n_nodes(g) && match_gated_residual(mat, next, ggml_graph_node(g, first + 2), op))
+ return 3;
+#endif
if (next->op == GGML_OP_CONCAT && next->src[1] == mat && next->src[0] != mat && next->op_params[0] == 1 &&
next->type == GGML_TYPE_F32 && mat->ne[2] == 1 && mat->ne[3] == 1 &&
next->ne[2] == 1 && next->ne[3] == 1 && ggml_is_contiguous(next) && ggml_is_contiguous(next->src[0]))
@@ -470,7 +531,8 @@ namespace SD_CUDA_NAMESPACE {
ggml_nelements(next) > int64_t(INT_MAX) * 256)
return 0;
const auto activation = ggml_get_glu_op(next);
- if (activation == GGML_GLU_OP_SITU && next->op_params[1] != 0) {
+ const bool situ = activation == GGML_GLU_OP_SITU && next->op_params[1] != 0;
+ if (situ) {
#ifdef SD_CUDA_I8_CUTLASS
if (!next->data && !next->buffer && !next->view_src && cutlass_compatible(op, next)) {
size_t workspace = 0;
@@ -480,9 +542,11 @@ namespace SD_CUDA_NAMESPACE {
return supported ? 2 : 0;
}
#endif
+#if !defined(SD_CUDA_VENDOR_HIP)
return 0;
+#endif
}
- if (activation != GGML_GLU_OP_SWIGLU || next->op_params[1] != 0)
+ if (!situ && (activation != GGML_GLU_OP_SWIGLU || next->op_params[1] != 0))
return 0;
if (first + 2 < ggml_graph_n_nodes(g)) {
const auto* quant = ggml_graph_node(g, first + 2);
@@ -504,7 +568,7 @@ namespace SD_CUDA_NAMESPACE {
return 3;
}
}
- return 2;
+ return situ ? 0 : 2;
}
struct Plan {
ggml_ext_gpu_device_info device;
@@ -529,7 +593,7 @@ namespace SD_CUDA_NAMESPACE {
*state = p;
auto device = ggml_backend_get_device(backend);
p->device = cuda_device_info(device);
- p->kind = !mat ? QUANTIZE : count == 3 ? SWIGLU_QUANTIZE
+ p->kind = !mat ? QUANTIZE : count == 3 ? (dst->op == GGML_OP_ADD ? GATED_RESIDUAL : SWIGLU_QUANTIZE)
: count == 1 ? MAT
: dst->op == GGML_OP_CONCAT ? CONCAT
: GATED;
@@ -538,7 +602,7 @@ namespace SD_CUDA_NAMESPACE {
p->rows = mat ? ggml_nrows(nodes[0]) : ggml_nrows(dst->src[0]);
p->padded = (p->rows + 3) / 4 * 4;
p->group = mat ? op.group_size : 0;
- p->situ = p->kind == GATED && ggml_get_glu_op(dst) == GGML_GLU_OP_SITU;
+ p->situ = (p->kind == GATED || p->kind == SWIGLU_QUANTIZE) && ggml_get_glu_op(nodes[1]) == GGML_GLU_OP_SITU;
if (!mat)
sd_ggml_convrot_describe(dst, &p->group);
size_t total = 0;
@@ -741,10 +805,10 @@ namespace SD_CUDA_NAMESPACE {
}
dst[i] = situ ? sd_convrot::situ(up, gate) : __fmul_rn(gate / (1.f + expf(-gate)), up);
}
- template <bool Swiglu, bool Dequantize, bool Rotated>
+ template <bool Swiglu, bool Dequantize, bool Rotated, bool Situ = false>
static CUDAError quantize_epilogue(const Plan& p, const float* src, const float* xs, const float* ws, int8_t* output, CUDAStream stream) {
const int64_t hidden = p.n / 2;
- auto kernel = sd_convrot::quantize<Swiglu, Dequantize, false, false, Rotated>;
+ auto kernel = sd_convrot::quantize<Swiglu, Dequantize, false, false, Rotated, float, false, false, false, Situ>;
static std::mutex mutex;
static std::unordered_map<int, bool> initialized;
{
@@ -761,12 +825,17 @@ namespace SD_CUDA_NAMESPACE {
if (error != cuda_success)
return error;
}
- const int threads = Rotated && !Swiglu && !Dequantize ? (hidden <= 1024 ? 128 : hidden <= 8192 ? 256
- : 1024)
- : 256;
- kernel<<<unsigned(p.rows), threads, size_t(hidden) * sizeof(float), stream>>>(src, output,
- reinterpret_cast<float*>(output + hidden * p.padded), hidden, p.rows, hidden,
- Swiglu ? src + hidden : nullptr, Swiglu ? p.n : hidden, Swiglu ? p.n : hidden, xs, ws, nullptr, 0.f, 1.f, 0.f, nullptr, nullptr);
+ int threads = Situ && hidden >= 8192 ? 512 : 256;
+ if constexpr (Rotated && !Swiglu && !Dequantize) {
+ threads = hidden <= 1024 ? 128 : hidden <= 8192 ? 256
+ : 1024;
+ }
+ const auto* gate = Situ ? src + hidden : src;
+ const auto* up = Swiglu ? src + (Situ ? 0 : hidden) : nullptr;
+ const auto stride = Swiglu ? p.n : hidden;
+ kernel<<<unsigned(p.rows), threads, size_t(hidden) * sizeof(float), stream>>>(
+ gate, output, reinterpret_cast<float*>(output + hidden * p.padded), hidden, p.rows, hidden,
+ up, stride, stride, xs, ws, nullptr, 0.f, 1.f, 0.f, nullptr, nullptr);
return cuda_last_error();
}
static ggml_status execute(const ggml_ext_execution_context* execution, ggml_tensor* const* nodes, int count, void* state, void*) {
@@ -780,14 +849,15 @@ namespace SD_CUDA_NAMESPACE {
auto scratch_at = [scratch](size_t offset) { return scratch ? scratch + offset : nullptr; };
auto* dst = nodes[count - 1];
SDI8Operation op;
- const bool mat = sd_ggml_i8_describe(nodes[0], &op);
- bool alias = false;
+ const bool mat = sd_ggml_i8_describe(nodes[0], &op);
+ const auto* residual = p.kind == GATED_RESIDUAL ? residual_input(nodes[1], dst) : nullptr;
+ bool alias = false;
for (int i = 0; i < count; ++i)
for (auto* input : nodes[i]->src) {
bool internal = false;
for (int j = 0; j < count; ++j)
internal |= input == nodes[j];
- if (!internal)
+ if (!internal && !(residual && i == count - 1 && input == residual && dst->data == residual->data))
alias |= overlaps(dst, input);
}
// CUTLASS quantization epilogues read only the workspace intermediate after GEMM completes.
@@ -850,19 +920,16 @@ namespace SD_CUDA_NAMESPACE {
#if defined(SD_CUDA_VENDOR_MUSA)
return GGML_STATUS_FAILED;
#else
- const int32_t alpha = 1, beta = 0;
#if defined(SD_CUDA_VENDOR_HIP)
-#if HIP_VERSION >= 60500000
- constexpr auto compute_type = HIPBLAS_COMPUTE_32I;
+ const HipI8Gemm gemm{int(p.n), int(p.padded), int(p.k), op.weight->data, quant, accum,
+ static_cast<hipblasHandle_t>(handle), stream};
+#ifdef SD_HIP_I8_LT
+ if (!hip_i8_gemm_lt(context.device, gemm) && gemm.run_hipblas() != HIPBLAS_STATUS_SUCCESS)
#else
- constexpr auto compute_type = HIPBLAS_R_32I;
+ if (gemm.run_hipblas() != HIPBLAS_STATUS_SUCCESS)
#endif
- const auto status = hipblasGemmEx(static_cast<hipblasHandle_t>(handle), HIPBLAS_OP_T, HIPBLAS_OP_N,
- int(p.n), int(p.padded), int(p.k), &alpha,
- op.weight->data, HIPBLAS_R_8I, int(p.k), quant, HIPBLAS_R_8I, int(p.k),
- &beta, accum, HIPBLAS_R_32I, int(p.n), compute_type, HIPBLAS_GEMM_DEFAULT);
- if (status != HIPBLAS_STATUS_SUCCESS)
#else
+ const int32_t alpha = 1, beta = 0;
const auto status = cublasGemmEx(static_cast<cublasHandle_t>(handle), CUBLAS_OP_T, CUBLAS_OP_N,
int(p.n), int(p.padded), int(p.k), &alpha,
op.weight->data, CUDA_R_8I, int(p.k), quant, CUDA_R_8I, int(p.k),
@@ -878,6 +945,10 @@ namespace SD_CUDA_NAMESPACE {
error = quantize_epilogue<false, false, true>(p, values, nullptr, nullptr, static_cast<int8_t*>(output), stream);
else if (p.gemm)
error = quantize_epilogue<true, false, false>(p, values, nullptr, nullptr, static_cast<int8_t*>(output), stream);
+#if defined(SD_CUDA_VENDOR_HIP)
+ else if (p.situ)
+ error = quantize_epilogue<true, true, false, true>(p, values, scales, ws, static_cast<int8_t*>(output), stream);
+#endif
else
error = quantize_epilogue<true, true, false>(p, values, scales, ws, static_cast<int8_t*>(output), stream);
} else if (p.kind == GATED && !p.fused_activation) {
@@ -889,7 +960,15 @@ namespace SD_CUDA_NAMESPACE {
error = cuda_last_error();
} else if (!p.gemm) {
const unsigned gy = unsigned(std::min<int64_t>(p.rows, 65535)), gz = unsigned((p.rows + 65534) / 65535);
- if ((reinterpret_cast<uintptr_t>(ws) | reinterpret_cast<uintptr_t>(bias) | reinterpret_cast<uintptr_t>(gemm_output)) % 16 == 0) {
+#if defined(SD_CUDA_VENDOR_HIP)
+ if (p.kind == GATED_RESIDUAL) {
+ dequantize_gated_residual_vec4<<<dim3(unsigned((p.n + 1023) / 1024), gy, gz), 256, 0, stream>>>(
+ reinterpret_cast<const int4*>(accum), scales, reinterpret_cast<const float4*>(ws),
+ static_cast<const float4*>(nodes[1]->src[1]->data), static_cast<const float4*>(residual->data),
+ reinterpret_cast<float4*>(gemm_output), p.n / 4, p.rows);
+ } else
+#endif
+ if ((reinterpret_cast<uintptr_t>(ws) | reinterpret_cast<uintptr_t>(bias) | reinterpret_cast<uintptr_t>(gemm_output)) % 16 == 0) {
dequantize_i32_rows_vec4_cuda<<<dim3(unsigned((p.n + 1023) / 1024), gy, gz), 256, 0, stream>>>(
reinterpret_cast<const int4*>(accum), scales, reinterpret_cast<const float4*>(ws), reinterpret_cast<const float4*>(bias), reinterpret_cast<float4*>(gemm_output), p.n / 4, p.rows);
} else {
diff --git a/src/core/cuda/registration.cu b/src/core/cuda/registration.cu
index d92be44..2755264 100644
--- a/src/core/cuda/registration.cu
+++ b/src/core/cuda/registration.cu
@@ -16,6 +16,9 @@ namespace SD_CUDA_NAMESPACE {
void register_modulate(ggml_backend_dev_t device);
void register_bf16(ggml_backend_dev_t device);
void register_rms_affine(ggml_backend_dev_t device);
+#if defined(SD_CUDA_VENDOR_HIP)
+ void register_wmma_attention(ggml_backend_dev_t device);
+#endif
#if !defined(SD_CUDA_VENDOR_HIP) && !defined(SD_CUDA_VENDOR_MUSA)
void register_sage_attention(ggml_backend_dev_t device);
void register_sol_attention(ggml_backend_dev_t device);
@@ -32,6 +35,9 @@ namespace SD_CUDA_NAMESPACE {
register_rope(device);
register_bf16(device);
register_rms_affine(device);
+#if defined(SD_CUDA_VENDOR_HIP)
+ register_wmma_attention(device);
+#endif
#if !defined(SD_CUDA_VENDOR_HIP) && !defined(SD_CUDA_VENDOR_MUSA)
register_sage_attention(device);
register_sol_attention(device);