[PR #264] Compile for sm_86 and keep its native int<->float conversions #270

New issue

Open

opened 2026-10-26 22:22:18 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-26 22:22:18 +00:00

Owner

📋 Pull Request Information

Original PR: https://github.com/Comfy-Org/comfy-kitchen/pull/264
Author: @Many0therFunctions
Created: 10/20/2026
Status: 🔄 Open

Base: main ← Head: sm86-native-conversions

📝 Commits (1)

86dfb29 Compile for sm_86 and keep its native int<->float conversions

📊 Changes

3 files changed (+13 additions, -9 deletions)

View changed files

📝 comfy_kitchen/backends/cuda/sage_attention/math.cuh (+5 -3)
📝 comfy_kitchen/backends/cuda/sage_attention/mma.cuh (+6 -4)
📝 setup.py (+2 -2)

📄 Description

What

Two changes for GA10x cards (RTX 3060-3090, A6000, A40):

setup.py: add 86-real to both default arch lists. The wheels have never carried sm_86 SASS; an sm_86 card runs the sm_80 cubin, which compiles int<->float conversions to I2F / F2I on the XU pipe. The sm_86 cubin uses I2FP / F2IP instead, which Nsight Compute places on the ALU pipe (one value per I2FP; F2IP rounds, saturates and packs two values per instruction), with the XU instruction count unchanged, as on sm_89.
sage_attention/math.cuh, mma.cuh: the arithmetic conversion forms from #234 / #235 (i2f_exact, pack_u8x4_exact, pack_u8x4_exact_nonneg) are gated __CUDA_ARCH__ < 860 instead of < 890. They were measured on the sm_80 cubin (my #235 numbers say so: "A6000, sm_80 cubin from the default arch manifest"); on native sm_86 code they lose to the plain conversions everywhere.

Nothing changes for sm_75 / sm_80 (still the arithmetic forms) or sm_89+ (still plain).

Why

Throughput of the conversion forms on an RTX A6000, elements per SM per clock (a 6-form microbenchmark, same source built -arch=sm_80 and -arch=sm_86, both run on the same card; source below):

form	sm_80 cubin	sm_86 cubin
int32 -> fp32, `__int2float_rn` (I2F / I2FP)	15.6	28.3
int32 -> fp32, `x + 0x4B400000` then `- 1.5*2^23`	29.3	24.3
4x fp32 -> packed u8, `cvt.rni` + `cvt.pack.sat` (F2I+I2IP / F2IP)	14.3	23.8
4x fp32 -> packed u8, `+2^23` + 3 PRMT	19.8	19.3
4x fp32 -> packed u8, `fminf 255` + `+2^23` + 3 PRMT (current `_nonneg`)	15.5	15.2

Nsight Compute on the same microbenchmark: on the sm_80 cubin the plain int->float adds one XU instruction per element; on the sm_86 cubin it adds one ALU instruction per element and no XU instructions. The arithmetic form adds one ALU (IADD3) and one FMA (FADD) per element on both.

Static SASS of the D128 128-row unmasked qk_int_sv_i8_attn_kernel (sm_86): v0.2.39 4144 instructions, this PR 3232 (-22%); the 128-row key-mask kernel 18544 -> 13072. No spills at D128; register counts within 2 of v0.2.39. Profiled at B1 H56 S53730 D128, the kernel executes 65.9 G warp instructions instead of 88.6 G, in 710 M cycles instead of 731 M, with the XU instruction count identical (5.22 G: the ex2 and rcp only) and the ALU pipe at 19% (v0.2.39: 24%).

Results

A6000 (sm_86, 300 W cap), both builds 86-real, sustained runs (1-3 s per arm, 3 rounds), interleaved A/B/B/A in separate processes, NVML-sampled, ck.int8_attention bf16 in / bf16 out, D128:

shape	path	v0.2.39	this PR	delta
H56 S53730	unmasked, 128-row tile	448.0 ms (1680 MHz, 286 W)	449.0 ms (1621 MHz, 290 W)	+0.2% (tie; see note)
H32 S16384	unmasked, 64-row tile	25.89 ms	24.30 ms	-6.1%
H32 S32768	unmasked, 64-row tile	99.80 ms	94.22 ms	-5.6%
H32 q16384 kv32768	prepared key mask	54.90 ms	51.08 ms	-7.0%
H56 S53730	prepared key mask	509.5 ms	462.6 ms	-9.2%

Note on the 128-row row: this PR needs 1.4% fewer cycles there but runs at a lower clock on this power-capped card. ncu shows the leaner kernel's L2 hit rate dropping from 98% to 60-80% (the co-resident CTAs stream the same K/V and only share it in L2 while they stay in step), so DRAM traffic rises from 2% to 26-30% of peak and the limiter takes the difference out of the clock. That is specific to a 300 W card; a 3080/3090 with more power headroom should see the cycle saving as time.

Output is bit-identical to v0.2.39 (and to the pre-#234 kernel) on D64 / D128 / D256, a partial tile and both masked paths. tests/test_int8_attention.py and tests/test_sol_attn.py pass on the sm_86 build (the two long_sequence_and_partial_tile[scale<=0] cases fail as before: torch's reference returns NaN there).

Without 86-real the gate change is inert for wheel users, who all run the sm_80 cubin; with it, the current kernels alone gain 2.4% at 53k tokens unmasked (sm_86 cubin vs sm_80 cubin of v0.2.39, same source). An sm_86 cubin adds about 4 MB to the Linux wheel.

Checking on a 3080 / 3090

Any GA10x card shows the same split: build once with COMFY_CUDA_ARCHS="80-real" and once with "86-real", cuobjdump -sass the two _C.abi3.so and grep the attention kernel for I2F vs I2FP; then time any int8 attention shape under 40k keys on both builds of this branch.

conv_throughput.cu (the microbenchmark; nvcc -O3 -arch=sm_86 conv_throughput.cu -o conv86 && ./conv86, likewise sm_80)

// Throughput of int32 -> fp32 and fp32 -> packed-u8 conversion forms on the current GPU.
// Each form runs in a throughput-bound loop (16 independent chains per thread, all blocks
// co-resident); elements per SM per clock from clock64() per block.
#include <cuda_runtime.h>
#include <cstdint>
#include <cstdio>

namespace {
constexpr int kElems = 16, kIters = 8192, kThreads = 256;
enum Form : int { kNone, kPlainI2F, kArithI2F, kPlainPack, kArithPack, kArithPackClamp };

const char* name(int form) {
  switch (form) {
    case kNone: return "loop overhead only (reinterpret, no conversion)";
    case kPlainI2F: return "int->float: __int2float_rn";
    case kArithI2F: return "int->float: x + 0x4B400000, - 1.5*2^23 (i2f_exact)";
    case kPlainPack: return "4x float->u8 pack: cvt.rni + cvt.pack.sat (pack_u8x4)";
    case kArithPack: return "4x float->u8 pack: +2^23, 3 PRMT (bounded)";
    default: return "4x float->u8 pack: fminf 255, +2^23, 3 PRMT (nonneg)";
  }
}

__device__ __forceinline__ uint32_t pack_plain(float a, float b, float c, float d) {
  uint32_t qa, qb, qc, qd, qdc, packed;
  asm volatile("cvt.rni.s32.f32 %0, %1;" : "=r"(qa) : "f"(a));
  asm volatile("cvt.rni.s32.f32 %0, %1;" : "=r"(qb) : "f"(b));
  asm volatile("cvt.rni.s32.f32 %0, %1;" : "=r"(qc) : "f"(c));
  asm volatile("cvt.rni.s32.f32 %0, %1;" : "=r"(qd) : "f"(d));
  asm volatile("cvt.pack.sat.u8.s32.b32 %0, %1, %2, 0;" : "=r"(qdc) : "r"(qd), "r"(qc));
  asm volatile("cvt.pack.sat.u8.s32.b32 %0, %1, %2, %3;" : "=r"(packed) : "r"(qb), "r"(qa), "r"(qdc));
  return packed;
}

template <bool clamp>
__device__ __forceinline__ uint32_t pack_arith(float a, float b, float c, float d) {
  if (clamp) { a = fminf(a, 255.f); b = fminf(b, 255.f); c = fminf(c, 255.f); d = fminf(d, 255.f); }
  const uint32_t qa = __float_as_uint(a + 8388608.0f), qb = __float_as_uint(b + 8388608.0f);
  const uint32_t qc = __float_as_uint(c + 8388608.0f), qd = __float_as_uint(d + 8388608.0f);
  return __byte_perm(__byte_perm(qa, qb, 0x0040), __byte_perm(qc, qd, 0x0040), 0x5410);
}

template <int form>
__global__ void bench(const int32_t* in, float* out, long long* cycles, float scale) {
  int32_t x[kElems];
  float acc[kElems];
  uint32_t packed_acc = 0;
#pragma unroll
  for (int j = 0; j < kElems; ++j) {
    x[j] = in[(threadIdx.x * kElems + j) & 4095] & 0xFF;
    acc[j] = 0.f;
  }
  __syncthreads();
  const long long t0 = clock64();
#pragma unroll 2
  for (int i = 0; i < kIters; ++i) {
    if constexpr (form <= kArithI2F) {
#pragma unroll
      for (int j = 0; j < kElems; ++j) {
        const float f = form == kNone ? __int_as_float(x[j])
                        : form == kPlainI2F ? __int2float_rn(x[j])
                                            : __int_as_float(x[j] + 0x4B400000) - 12582912.0f;
        acc[j] = fmaf(f, scale, acc[j]);
        x[j] ^= i & 0x7F;
      }
    } else {
#pragma unroll
      for (int j = 0; j < kElems; j += 4) {
        const float a = __int_as_float(0x43000000 | (x[j] & 0x7F) << 16);  // 128..255
        const float b = __int_as_float(0x43000000 | (x[j + 1] & 0x7F) << 16);
        const float c = __int_as_float(0x43000000 | (x[j + 2] & 0x7F) << 16);
        const float d = __int_as_float(0x43000000 | (x[j + 3] & 0x7F) << 16);
        const uint32_t p = form == kPlainPack ? pack_plain(a, b, c, d)
                           : form == kArithPack ? pack_arith<false>(a, b, c, d)
                                                : pack_arith<true>(a, b, c, d);
        packed_acc = __dp4a(p, 0x01010101u, packed_acc);
#pragma unroll
        for (int t = 0; t < 4; ++t) x[j + t] ^= i & 0x7F;
      }
    }
  }
  const long long t1 = clock64();
  float s = static_cast<float>(packed_acc);
#pragma unroll
  for (int j = 0; j < kElems; ++j) s += acc[j];
  out[blockIdx.x * blockDim.x + threadIdx.x] = s;
  if (threadIdx.x == 0) cycles[blockIdx.x] = t1 - t0;
}

template <int form>
void run(int sms, int blocks_per_sm, const int32_t* in, float* out, long long* cycles_d) {
  const int blocks = sms * blocks_per_sm;
  bench<form><<<blocks, kThreads>>>(in, out, cycles_d, 1.0001f);  // warm-up
  bench<form><<<blocks, kThreads>>>(in, out, cycles_d, 1.0001f);
  cudaDeviceSynchronize();
  long long cycles[4096];
  cudaMemcpy(cycles, cycles_d, blocks * sizeof(long long), cudaMemcpyDeviceToHost);
  long long worst = 0;
  for (int b = 0; b < blocks; ++b) worst = cycles[b] > worst ? cycles[b] : worst;
  const double elems_per_sm = double(blocks_per_sm) * kThreads * kIters * kElems;
  printf("  %-62s %6.1f elements/SM/clk\n", name(form), elems_per_sm / double(worst));
}
}  // namespace

int main() {
  cudaDeviceProp p{};
  cudaGetDeviceProperties(&p, 0);
  int occ = 0;
  cudaOccupancyMaxActiveBlocksPerMultiprocessor(&occ, bench<kArithPackClamp>, kThreads, 0);
  const int blocks_per_sm = occ < 4 ? occ : 4;
  printf("%s sm_%d%d, %d SMs, %d blocks x %d threads per SM\n", p.name, p.major, p.minor,
         p.multiProcessorCount, blocks_per_sm, kThreads);
  int32_t* in; float* out; long long* cycles;
  cudaMalloc(&in, 4096 * sizeof(int32_t));
  cudaMemset(in, 0x5A, 4096 * sizeof(int32_t));
  cudaMalloc(&out, p.multiProcessorCount * blocks_per_sm * kThreads * sizeof(float));
  cudaMalloc(&cycles, p.multiProcessorCount * blocks_per_sm * sizeof(long long));
  run<kNone>(p.multiProcessorCount, blocks_per_sm, in, out, cycles);
  run<kPlainI2F>(p.multiProcessorCount, blocks_per_sm, in, out, cycles);
  run<kArithI2F>(p.multiProcessorCount, blocks_per_sm, in, out, cycles);
  run<kPlainPack>(p.multiProcessorCount, blocks_per_sm, in, out, cycles);
  run<kArithPack>(p.multiProcessorCount, blocks_per_sm, in, out, cycles);
  run<kArithPackClamp>(p.multiProcessorCount, blocks_per_sm, in, out, cycles);
  const cudaError_t err = cudaGetLastError();
  if (err != cudaSuccess) { printf("CUDA error: %s\n", cudaGetErrorString(err)); return 1; }
  return 0;
}

_{🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.}

## 📋 Pull Request Information **Original PR:** https://github.com/Comfy-Org/comfy-kitchen/pull/264 **Author:** [@Many0therFunctions](https://github.com/Many0therFunctions) **Created:** 10/20/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `sm86-native-conversions` --- ### 📝 Commits (1) - [`86dfb29`](https://github.com/Comfy-Org/comfy-kitchen/commit/86dfb297c87b4aa00d1ff798903a5fe6d9769998) Compile for sm_86 and keep its native int<->float conversions ### 📊 Changes **3 files changed** (+13 additions, -9 deletions) <details> <summary>View changed files</summary> 📝 `comfy_kitchen/backends/cuda/sage_attention/math.cuh` (+5 -3) 📝 `comfy_kitchen/backends/cuda/sage_attention/mma.cuh` (+6 -4) 📝 `setup.py` (+2 -2) </details> ### 📄 Description ## What Two changes for GA10x cards (RTX 3060-3090, A6000, A40): 1. `setup.py`: add `86-real` to both default arch lists. The wheels have never carried sm_86 SASS; an sm_86 card runs the sm_80 cubin, which compiles int<->float conversions to `I2F` / `F2I` on the XU pipe. The sm_86 cubin uses `I2FP` / `F2IP` instead, which Nsight Compute places on the ALU pipe (one value per `I2FP`; `F2IP` rounds, saturates and packs two values per instruction), with the XU instruction count unchanged, as on sm_89. 2. `sage_attention/math.cuh`, `mma.cuh`: the arithmetic conversion forms from #234 / #235 (`i2f_exact`, `pack_u8x4_exact`, `pack_u8x4_exact_nonneg`) are gated `__CUDA_ARCH__ < 860` instead of `< 890`. They were measured on the sm_80 cubin (my #235 numbers say so: "A6000, sm_80 cubin from the default arch manifest"); on native sm_86 code they lose to the plain conversions everywhere. Nothing changes for sm_75 / sm_80 (still the arithmetic forms) or sm_89+ (still plain). ## Why Throughput of the conversion forms on an RTX A6000, elements per SM per clock (a 6-form microbenchmark, same source built `-arch=sm_80` and `-arch=sm_86`, both run on the same card; source below): | form | sm_80 cubin | sm_86 cubin | |---|---|---| | int32 -> fp32, `__int2float_rn` (I2F / I2FP) | 15.6 | **28.3** | | int32 -> fp32, `x + 0x4B400000` then `- 1.5*2^23` | **29.3** | 24.3 | | 4x fp32 -> packed u8, `cvt.rni` + `cvt.pack.sat` (F2I+I2IP / F2IP) | 14.3 | **23.8** | | 4x fp32 -> packed u8, `+2^23` + 3 PRMT | **19.8** | 19.3 | | 4x fp32 -> packed u8, `fminf 255` + `+2^23` + 3 PRMT (current `_nonneg`) | 15.5 | 15.2 | Nsight Compute on the same microbenchmark: on the sm_80 cubin the plain int->float adds one XU instruction per element; on the sm_86 cubin it adds one ALU instruction per element and no XU instructions. The arithmetic form adds one ALU (`IADD3`) and one FMA (`FADD`) per element on both. Static SASS of the D128 128-row unmasked `qk_int_sv_i8_attn_kernel` (sm_86): v0.2.39 4144 instructions, this PR 3232 (-22%); the 128-row key-mask kernel 18544 -> 13072. No spills at D128; register counts within 2 of v0.2.39. Profiled at B1 H56 S53730 D128, the kernel executes 65.9 G warp instructions instead of 88.6 G, in 710 M cycles instead of 731 M, with the XU instruction count identical (5.22 G: the `ex2` and `rcp` only) and the ALU pipe at 19% (v0.2.39: 24%). ## Results A6000 (sm_86, 300 W cap), both builds `86-real`, sustained runs (1-3 s per arm, 3 rounds), interleaved A/B/B/A in separate processes, NVML-sampled, `ck.int8_attention` bf16 in / bf16 out, D128: | shape | path | v0.2.39 | this PR | delta | |---|---|---|---|---| | H56 S53730 | unmasked, 128-row tile | 448.0 ms (1680 MHz, 286 W) | 449.0 ms (1621 MHz, 290 W) | +0.2% (tie; see note) | | H32 S16384 | unmasked, 64-row tile | 25.89 ms | 24.30 ms | **-6.1%** | | H32 S32768 | unmasked, 64-row tile | 99.80 ms | 94.22 ms | **-5.6%** | | H32 q16384 kv32768 | prepared key mask | 54.90 ms | 51.08 ms | **-7.0%** | | H56 S53730 | prepared key mask | 509.5 ms | 462.6 ms | **-9.2%** | Note on the 128-row row: this PR needs 1.4% fewer cycles there but runs at a lower clock on this power-capped card. ncu shows the leaner kernel's L2 hit rate dropping from 98% to 60-80% (the co-resident CTAs stream the same K/V and only share it in L2 while they stay in step), so DRAM traffic rises from 2% to 26-30% of peak and the limiter takes the difference out of the clock. That is specific to a 300 W card; a 3080/3090 with more power headroom should see the cycle saving as time. Output is bit-identical to v0.2.39 (and to the pre-#234 kernel) on D64 / D128 / D256, a partial tile and both masked paths. `tests/test_int8_attention.py` and `tests/test_sol_attn.py` pass on the sm_86 build (the two `long_sequence_and_partial_tile[scale<=0]` cases fail as before: torch's reference returns NaN there). Without `86-real` the gate change is inert for wheel users, who all run the sm_80 cubin; with it, the current kernels alone gain 2.4% at 53k tokens unmasked (sm_86 cubin vs sm_80 cubin of v0.2.39, same source). An sm_86 cubin adds about 4 MB to the Linux wheel. ## Checking on a 3080 / 3090 Any GA10x card shows the same split: build once with `COMFY_CUDA_ARCHS="80-real"` and once with `"86-real"`, `cuobjdump -sass` the two `_C.abi3.so` and grep the attention kernel for `I2F ` vs `I2FP`; then time any int8 attention shape under 40k keys on both builds of this branch. <details> <summary>conv_throughput.cu (the microbenchmark; <code>nvcc -O3 -arch=sm_86 conv_throughput.cu -o conv86 && ./conv86</code>, likewise sm_80)</summary> ```cpp // Throughput of int32 -> fp32 and fp32 -> packed-u8 conversion forms on the current GPU. // Each form runs in a throughput-bound loop (16 independent chains per thread, all blocks // co-resident); elements per SM per clock from clock64() per block. #include <cuda_runtime.h> #include <cstdint> #include <cstdio> namespace { constexpr int kElems = 16, kIters = 8192, kThreads = 256; enum Form : int { kNone, kPlainI2F, kArithI2F, kPlainPack, kArithPack, kArithPackClamp }; const char* name(int form) { switch (form) { case kNone: return "loop overhead only (reinterpret, no conversion)"; case kPlainI2F: return "int->float: __int2float_rn"; case kArithI2F: return "int->float: x + 0x4B400000, - 1.5*2^23 (i2f_exact)"; case kPlainPack: return "4x float->u8 pack: cvt.rni + cvt.pack.sat (pack_u8x4)"; case kArithPack: return "4x float->u8 pack: +2^23, 3 PRMT (bounded)"; default: return "4x float->u8 pack: fminf 255, +2^23, 3 PRMT (nonneg)"; } } __device__ __forceinline__ uint32_t pack_plain(float a, float b, float c, float d) { uint32_t qa, qb, qc, qd, qdc, packed; asm volatile("cvt.rni.s32.f32 %0, %1;" : "=r"(qa) : "f"(a)); asm volatile("cvt.rni.s32.f32 %0, %1;" : "=r"(qb) : "f"(b)); asm volatile("cvt.rni.s32.f32 %0, %1;" : "=r"(qc) : "f"(c)); asm volatile("cvt.rni.s32.f32 %0, %1;" : "=r"(qd) : "f"(d)); asm volatile("cvt.pack.sat.u8.s32.b32 %0, %1, %2, 0;" : "=r"(qdc) : "r"(qd), "r"(qc)); asm volatile("cvt.pack.sat.u8.s32.b32 %0, %1, %2, %3;" : "=r"(packed) : "r"(qb), "r"(qa), "r"(qdc)); return packed; } template <bool clamp> __device__ __forceinline__ uint32_t pack_arith(float a, float b, float c, float d) { if (clamp) { a = fminf(a, 255.f); b = fminf(b, 255.f); c = fminf(c, 255.f); d = fminf(d, 255.f); } const uint32_t qa = __float_as_uint(a + 8388608.0f), qb = __float_as_uint(b + 8388608.0f); const uint32_t qc = __float_as_uint(c + 8388608.0f), qd = __float_as_uint(d + 8388608.0f); return __byte_perm(__byte_perm(qa, qb, 0x0040), __byte_perm(qc, qd, 0x0040), 0x5410); } template <int form> __global__ void bench(const int32_t* in, float* out, long long* cycles, float scale) { int32_t x[kElems]; float acc[kElems]; uint32_t packed_acc = 0; #pragma unroll for (int j = 0; j < kElems; ++j) { x[j] = in[(threadIdx.x * kElems + j) & 4095] & 0xFF; acc[j] = 0.f; } __syncthreads(); const long long t0 = clock64(); #pragma unroll 2 for (int i = 0; i < kIters; ++i) { if constexpr (form <= kArithI2F) { #pragma unroll for (int j = 0; j < kElems; ++j) { const float f = form == kNone ? __int_as_float(x[j]) : form == kPlainI2F ? __int2float_rn(x[j]) : __int_as_float(x[j] + 0x4B400000) - 12582912.0f; acc[j] = fmaf(f, scale, acc[j]); x[j] ^= i & 0x7F; } } else { #pragma unroll for (int j = 0; j < kElems; j += 4) { const float a = __int_as_float(0x43000000 | (x[j] & 0x7F) << 16); // 128..255 const float b = __int_as_float(0x43000000 | (x[j + 1] & 0x7F) << 16); const float c = __int_as_float(0x43000000 | (x[j + 2] & 0x7F) << 16); const float d = __int_as_float(0x43000000 | (x[j + 3] & 0x7F) << 16); const uint32_t p = form == kPlainPack ? pack_plain(a, b, c, d) : form == kArithPack ? pack_arith<false>(a, b, c, d) : pack_arith<true>(a, b, c, d); packed_acc = __dp4a(p, 0x01010101u, packed_acc); #pragma unroll for (int t = 0; t < 4; ++t) x[j + t] ^= i & 0x7F; } } } const long long t1 = clock64(); float s = static_cast<float>(packed_acc); #pragma unroll for (int j = 0; j < kElems; ++j) s += acc[j]; out[blockIdx.x * blockDim.x + threadIdx.x] = s; if (threadIdx.x == 0) cycles[blockIdx.x] = t1 - t0; } template <int form> void run(int sms, int blocks_per_sm, const int32_t* in, float* out, long long* cycles_d) { const int blocks = sms * blocks_per_sm; bench<form><<<blocks, kThreads>>>(in, out, cycles_d, 1.0001f); // warm-up bench<form><<<blocks, kThreads>>>(in, out, cycles_d, 1.0001f); cudaDeviceSynchronize(); long long cycles[4096]; cudaMemcpy(cycles, cycles_d, blocks * sizeof(long long), cudaMemcpyDeviceToHost); long long worst = 0; for (int b = 0; b < blocks; ++b) worst = cycles[b] > worst ? cycles[b] : worst; const double elems_per_sm = double(blocks_per_sm) * kThreads * kIters * kElems; printf(" %-62s %6.1f elements/SM/clk\n", name(form), elems_per_sm / double(worst)); } } // namespace int main() { cudaDeviceProp p{}; cudaGetDeviceProperties(&p, 0); int occ = 0; cudaOccupancyMaxActiveBlocksPerMultiprocessor(&occ, bench<kArithPackClamp>, kThreads, 0); const int blocks_per_sm = occ < 4 ? occ : 4; printf("%s sm_%d%d, %d SMs, %d blocks x %d threads per SM\n", p.name, p.major, p.minor, p.multiProcessorCount, blocks_per_sm, kThreads); int32_t* in; float* out; long long* cycles; cudaMalloc(&in, 4096 * sizeof(int32_t)); cudaMemset(in, 0x5A, 4096 * sizeof(int32_t)); cudaMalloc(&out, p.multiProcessorCount * blocks_per_sm * kThreads * sizeof(float)); cudaMalloc(&cycles, p.multiProcessorCount * blocks_per_sm * sizeof(long long)); run<kNone>(p.multiProcessorCount, blocks_per_sm, in, out, cycles); run<kPlainI2F>(p.multiProcessorCount, blocks_per_sm, in, out, cycles); run<kArithI2F>(p.multiProcessorCount, blocks_per_sm, in, out, cycles); run<kPlainPack>(p.multiProcessorCount, blocks_per_sm, in, out, cycles); run<kArithPack>(p.multiProcessorCount, blocks_per_sm, in, out, cycles); run<kArithPackClamp>(p.multiProcessorCount, blocks_per_sm, in, out, cycles); const cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { printf("CUDA error: %s\n", cudaGetErrorString(err)); return 1; } return 0; } ``` </details> --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>

bsakaguchi added the

pull-request

label

2026-10-26 22:22:18 +00:00

bsakaguchi referenced this issue

2026-10-26 22:22:20 +00:00

[PR #270] Add Metal RoPE and AdaLN kernels for Apple GPUs (mps backend), based on main @ f8b7e97 (v0.2.40) #273

bsakaguchi referenced this issue

2026-10-26 22:22:23 +00:00

[PR #275] Add ck.input_act with Metal swiglu/situ for Apple GPUs (depends on #270), based on main @ f8b7e97 (v0.2.40) #277