incorrect
17_FP8ScaledAttention.py
started
2026-05-11T22:16:06Z
correctness issue:
timingkernelbench.timing
kernel mean / median / std
—
ref mean / median / std
42.90 / — / 1.80 μs
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import math
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Scaled low-precision attention with explicit dequantization semantics.
The inputs are quantized tensors plus per-head scales. This mirrors the data
path of FP8-style scaled matmul attention even though the reference computes
everything through explicit dequantization in PyTorch.
"""
def __init__(self):
super().__init__()
def forward(
self,
q_q: torch.Tensor,
k_q: torch.Tensor,
v_q: torch.Tensor,
q_scale: torch.Tensor,
k_scale: torch.Tensor,
v_scale: torch.Tensor,
) -> torch.Tensor:
q = q_q.float() * q_scale.view(1, 1, -1, 1)
k = k_q.float() * k_scale.view(1, 1, -1, 1)
v = v_q.float() * v_scale.view(1, 1, -1, 1)
batch_size, seq_len, num_heads, head_dim = q.shape
out = torch.zeros_like(q)
scale = 1.0 / math.sqrt(head_dim)
for b in range(batch_size):
for h in range(num_heads):
for t in range(seq_len):
q_t = q[b, t, h]
k_hist = k[b, : t + 1, h]
v_hist = v[b, : t + 1, h]
scores = (k_hist * q_t.unsqueeze(0)).sum(dim=-1) * scale
attn = torch.softmax(scores, dim=0)
out[b, t, h] = (attn.unsqueeze(-1) * v_hist).sum(dim=0)
return out
batch_size = 2
seq_len = 72
num_heads = 4
head_dim = 32
def _quantize_per_head(x: torch.Tensor):
scale = x.abs().amax(dim=(0, 1, 3)).clamp(min=1e-4) / 127.0
q = torch.clamp(torch.round(x / scale.view(1, 1, -1, 1)), -127, 127).to(torch.int8)
return q, scale.to(torch.float32)
def get_inputs():
p = popcorn_pri
q = torch.randn(p.jitter_int(batch_size), p.jitter_int(seq_len), p.jitter_int(num_heads), p.jitter_int(head_dim, align=8), dtype=torch.float32)
k = torch.randn(p.jitter_int(batch_size), p.jitter_int(seq_len), p.jitter_int(num_heads), p.jitter_int(head_dim, align=8), dtype=torch.float32)
v = torch.randn(p.jitter_int(batch_size), p.jitter_int(seq_len), p.jitter_int(num_heads), p.jitter_int(head_dim, align=8), dtype=torch.float32)
q_q, q_scale = _quantize_per_head(q)
k_q, k_scale = _quantize_per_head(k)
v_q, v_scale = _quantize_per_head(v)
return [q_q, k_q, v_q, q_scale, k_scale, v_scale]
def get_init_inputs():
return []
turn 0244.2s turn · 16.7s LLM · 227.4s tools · tokens in=2,393 · out=1,081 · total=4,111
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The reference implementation is a scaled low-precision attention with explicit dequantization, but it's ...The task requires replacing the forward computation of a PyTorch model with a custom CUDA kernel. The original model performs scaled low-precision attention with explicit dequantization, but it's inefficient due to nested Python loops. I need to create a ModelNew class that mimics the original constructor arguments and forward signature, producing the same output: 8-bit quantized queries, keys, and values with their scales.
The computation involves dequantizing int8-quantized q, k, v tensors (shape [B, T, H, D]) using per-head scales, then performing causal self-attention.
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void causal_attention_kernel(\n const int8_t* __restrict__ q_q,\n const int8_t* __restrict__ k_q,\n const int8_t* __restrict__ v_q,\n const float* __restrict__ q_scale,\n const float* __restrict__ k_scale,\n const float* __restrict__ v_scale,\n float* __restrict__ out,\n int batch_size,\n int seq_len,\n int num_heads,\n int head_dim,\n float scale\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = batch_size * num_heads * seq_len;\n if (idx >= total) return;\n\n int t = idx % seq_len;\n int h = (idx / seq_len) % num_heads;\n int b = idx / (seq_len * num_heads);\n\n float q_scale_h = q_scale[h];\n float k_scale_h = k_scale[h];\n float v_scale_h = v_scale[h];\n\n float scores[72]; // max seq_len\n float max_score = -1e30f;\n for (int s = 0; s <= t; ++s) {\n float dot = 0.0f;\n for (int d = 0; d < head_dim; ++d) {\n float q_val = q_q[((b * seq_len + t) * num_heads + h) * head_dim + d] * q_scale_h;\n float k_val = k_q[((b * seq_len + s) * num_heads + h) * head_dim + d] * k_scale_h;\n dot += q_val * k_val;\n }\n scores[s] = dot * scale;\n if (scores[s] > max_score) max_score = scores[s];\n }\n\n float sum_exp = 0.0f;\n for (int s = 0; s <= t; ++s) {\n scores[s] = expf(scores[s] - max_score);\n sum_exp += scores[s];\n }\n float inv_sum = 1.0f / sum_exp;\n\n for (int d = 0; d < head_dim; ++d) {\n float val = 0.0f;\n for (int s = 0; s <= t; ++s) {\n float v_val = v_q[((b * seq_len + s) * num_heads + h) * head_dim + d] * v_scale_h;\n val += scores[s] * inv_sum * v_val;\n }\n out[((b * seq_len + t) * num_heads + h) * head_dim + d] = val;\n }\n}\n\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n) {\n auto batch_size = q_q.size(0);\n auto seq_len = q_q.size(1);\n auto num_heads = q_q.size(2);\n auto head_dim = q_q.size(3);\n\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(q_q.device());\n auto out = torch::zeros({batch_size, seq_len, num_heads, head_dim}, options);\n\n float scale = 1.0f / sqrtf((float)head_dim);\n\n int total = batch_size * num_heads * seq_len;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n causal_attention_kernel<<<blocks, threads>>>(\n q_q.data_ptr<int8_t>(),\n k_q.data_ptr<int8_t>(),\n v_q.data_ptr<int8_t>(),\n q_scale.data_ptr<float>(),\n k_scale.data_ptr<float>(),\n v_scale.data_ptr<float>(),\n out.data_ptr<float>(),\n batch_size,\n seq_len,\n num_heads,\n head_dim,\n scale\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"causal_attention\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"causal_attention_cuda\"],\n verbose=True,\n extra_cuda_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, q_q, k_q, v_q, q_scale, k_scale, v_scale):\n return self.module.causal_attention_cuda(q_q, k_q, v_q, q_scale, k_scale, v_scale)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 1273.4s turn · 6.8s LLM · 266.6s tools · tokens in=4,128 · out=1,082 · total=5,217
reasoning
Compiled successfully. Now check correctness.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void causal_attention_kernel(\n const int8_t* __restrict__ q_q,\n const int8_t* __restrict__ k_q,\n const int8_t* __restrict__ v_q,\n const float* __restrict__ q_scale,\n const float* __restrict__ k_scale,\n const float* __restrict__ v_scale,\n float* __restrict__ out,\n int batch_size,\n int seq_len,\n int num_heads,\n int head_dim,\n float scale\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = batch_size * num_heads * seq_len;\n if (idx >= total) return;\n\n int t = idx % seq_len;\n int h = (idx / seq_len) % num_heads;\n int b = idx / (seq_len * num_heads);\n\n float q_scale_h = q_scale[h];\n float k_scale_h = k_scale[h];\n float v_scale_h = v_scale[h];\n\n float scores[72]; // max seq_len\n float max_score = -1e30f;\n for (int s = 0; s <= t; ++s) {\n float dot = 0.0f;\n for (int d = 0; d < head_dim; ++d) {\n float q_val = q_q[((b * seq_len + t) * num_heads + h) * head_dim + d] * q_scale_h;\n float k_val = k_q[((b * seq_len + s) * num_heads + h) * head_dim + d] * k_scale_h;\n dot += q_val * k_val;\n }\n scores[s] = dot * scale;\n if (scores[s] > max_score) max_score = scores[s];\n }\n\n float sum_exp = 0.0f;\n for (int s = 0; s <= t; ++s) {\n scores[s] = expf(scores[s] - max_score);\n sum_exp += scores[s];\n }\n float inv_sum = 1.0f / sum_exp;\n\n for (int d = 0; d < head_dim; ++d) {\n float val = 0.0f;\n for (int s = 0; s <= t; ++s) {\n float v_val = v_q[((b * seq_len + s) * num_heads + h) * head_dim + d] * v_scale_h;\n val += scores[s] * inv_sum * v_val;\n }\n out[((b * seq_len + t) * num_heads + h) * head_dim + d] = val;\n }\n}\n\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n) {\n auto batch_size = q_q.size(0);\n auto seq_len = q_q.size(1);\n auto num_heads = q_q.size(2);\n auto head_dim = q_q.size(3);\n\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(q_q.device());\n auto out = torch::zeros({batch_size, seq_len, num_heads, head_dim}, options);\n\n float scale = 1.0f / sqrtf((float)head_dim);\n\n int total = batch_size * num_heads * seq_len;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n causal_attention_kernel<<<blocks, threads>>>(\n q_q.data_ptr<int8_t>(),\n k_q.data_ptr<int8_t>(),\n v_q.data_ptr<int8_t>(),\n q_scale.data_ptr<float>(),\n k_scale.data_ptr<float>(),\n v_scale.data_ptr<float>(),\n out.data_ptr<float>(),\n batch_size,\n seq_len,\n num_heads,\n head_dim,\n scale\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"causal_attention\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"causal_attention_cuda\"],\n verbose=True,\n extra_cuda_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, q_q, k_q, v_q, q_scale, k_scale, v_scale):\n return self.module.causal_attention_cuda(q_q, k_q, v_q, q_scale, k_scale, v_scale)\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: ? trials did not all match.
runtime_error: expected scalar type Char but found Float
turn 2248.2s turn · 11.1s LLM · 237.1s tools · tokens in=5,246 · out=1,156 · total=6,607
reasoning
Error: expected scalar type Char but found Float. The inputs q_q etc. are on CPU? No, probably need to move to CUDA or check types.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void causal_attention_kernel(\n const int8_t* __restrict__ q_q,\n const int8_t* __restrict__ k_q,\n const int8_t* __restrict__ v_q,\n const float* __restrict__ q_scale,\n const float* __restrict__ k_scale,\n const float* __restrict__ v_scale,\n float* __restrict__ out,\n int batch_size,\n int seq_len,\n int num_heads,\n int head_dim,\n float scale\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = batch_size * num_heads * seq_len;\n if (idx >= total) return;\n\n int t = idx % seq_len;\n int h = (idx / seq_len) % num_heads;\n int b = idx / (seq_len * num_heads);\n\n float q_scale_h = q_scale[h];\n float k_scale_h = k_scale[h];\n float v_scale_h = v_scale[h];\n\n float scores[72];\n float max_score = -1e30f;\n for (int s = 0; s <= t; ++s) {\n float dot = 0.0f;\n for (int d = 0; d < head_dim; ++d) {\n float q_val = static_cast<float>(q_q[((b * seq_len + t) * num_heads + h) * head_dim + d]) * q_scale_h;\n float k_val = static_cast<float>(k_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * k_scale_h;\n dot += q_val * k_val;\n }\n scores[s] = dot * scale;\n if (scores[s] > max_score) max_score = scores[s];\n }\n\n float sum_exp = 0.0f;\n for (int s = 0; s <= t; ++s) {\n scores[s] = expf(scores[s] - max_score);\n sum_exp += scores[s];\n }\n float inv_sum = 1.0f / sum_exp;\n\n for (int d = 0; d < head_dim; ++d) {\n float val = 0.0f;\n for (int s = 0; s <= t; ++s) {\n float v_val = static_cast<float>(v_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * v_scale_h;\n val += scores[s] * inv_sum * v_val;\n }\n out[((b * seq_len + t) * num_heads + h) * head_dim + d] = val;\n }\n}\n\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n) {\n auto batch_size = q_q.size(0);\n auto seq_len = q_q.size(1);\n auto num_heads = q_q.size(2);\n auto head_dim = q_q.size(3);\n\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(q_q.device());\n auto out = torch::zeros({batch_size, seq_len, num_heads, head_dim}, options);\n\n float scale = 1.0f / sqrtf((float)head_dim);\n\n int total = batch_size * num_heads * seq_len;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n causal_attention_kernel<<<blocks, threads>>>(\n q_q.data_ptr<int8_t>(),\n k_q.data_ptr<int8_t>(),\n v_q.data_ptr<int8_t>(),\n q_scale.data_ptr<float>(),\n k_scale.data_ptr<float>(),\n v_scale.data_ptr<float>(),\n out.data_ptr<float>(),\n batch_size,\n seq_len,\n num_heads,\n head_dim,\n scale\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"causal_attention\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"causal_attention_cuda\"],\n verbose=False,\n extra_cuda_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, q_q, k_q, v_q, q_scale, k_scale, v_scale):\n q_q = q_q.cuda().contiguous()\n k_q = k_q.cuda().contiguous()\n v_q = v_q.cuda().contiguous()\n q_scale = q_scale.cuda().contiguous()\n k_scale = k_scale.cuda().contiguous()\n v_scale = v_scale.cuda().contiguous()\n return self.module.causal_attention_cuda(q_q, k_q, v_q, q_scale, k_scale, v_scale).cpu()\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: ? trials did not all match.
runtime_error: expected scalar type Char but found Float
turn 3253.3s turn · 7.8s LLM · 245.5s tools · tokens in=6,636 · out=1,088 · total=7,845
reasoning
Still error. The error is probably because in the test, it expects the forward to handle CPU or something? No.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void causal_attention_kernel(\n const int8_t* __restrict__ q_q,\n const int8_t* __restrict__ k_q,\n const int8_t* __restrict__ v_q,\n const float* __restrict__ q_scale,\n const float* __restrict__ k_scale,\n const float* __restrict__ v_scale,\n float* __restrict__ out,\n int batch_size,\n int seq_len,\n int num_heads,\n int head_dim,\n float scale\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = batch_size * num_heads * seq_len;\n if (idx >= total) return;\n\n int t = idx % seq_len;\n int h = (idx / seq_len) % num_heads;\n int b = idx / (seq_len * num_heads);\n\n float q_scale_h = q_scale[h];\n float k_scale_h = k_scale[h];\n float v_scale_h = v_scale[h];\n\n float scores[72];\n float max_score = -1e30f;\n for (int s = 0; s <= t; ++s) {\n float dot = 0.0f;\n for (int d = 0; d < head_dim; ++d) {\n float q_val = static_cast<float>(q_q[((b * seq_len + t) * num_heads + h) * head_dim + d]) * q_scale_h;\n float k_val = static_cast<float>(k_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * k_scale_h;\n dot += q_val * k_val;\n }\n scores[s] = dot * scale;\n if (scores[s] > max_score) max_score = scores[s];\n }\n\n float sum_exp = 0.0f;\n for (int s = 0; s <= t; ++s) {\n scores[s] = expf(scores[s] - max_score);\n sum_exp += scores[s];\n }\n float inv_sum = 1.0f / sum_exp;\n\n for (int d = 0; d < head_dim; ++d) {\n float val = 0.0f;\n for (int s = 0; s <= t; ++s) {\n float v_val = static_cast<float>(v_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * v_scale_h;\n val += scores[s] * inv_sum * v_val;\n }\n out[((b * seq_len + t) * num_heads + h) * head_dim + d] = val;\n }\n}\n\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n) {\n auto batch_size = q_q.size(0);\n auto seq_len = q_q.size(1);\n auto num_heads = q_q.size(2);\n auto head_dim = q_q.size(3);\n\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(q_q.device());\n auto out = torch::zeros({batch_size, seq_len, num_heads, head_dim}, options);\n\n float scale = 1.0f / sqrtf((float)head_dim);\n\n int total = batch_size * num_heads * seq_len;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n causal_attention_kernel<<<blocks, threads>>>(\n q_q.data_ptr<int8_t>(),\n k_q.data_ptr<int8_t>(),\n v_q.data_ptr<int8_t>(),\n q_scale.data_ptr<float>(),\n k_scale.data_ptr<float>(),\n v_scale.data_ptr<float>(),\n out.data_ptr<float>(),\n batch_size,\n seq_len,\n num_heads,\n head_dim,\n scale\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"causal_attention\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"causal_attention_cuda\"],\n verbose=False,\n extra_cuda_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, q_q, k_q, v_q, q_scale, k_scale, v_scale):\n return self.module.causal_attention_cuda(q_q, k_q, v_q, q_scale, k_scale, v_scale)\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: ? trials did not all match.
runtime_error: expected scalar type Char but found Float
turn 4329.6s turn · 9.3s LLM · 320.4s tools · tokens in=7,874 · out=1,168 · total=9,222
reasoning
Still same. The problem is likely that the input q_q in test is float? No.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void causal_attention_kernel(\n const int8_t* __restrict__ q_q,\n const int8_t* __restrict__ k_q,\n const int8_t* __restrict__ v_q,\n const float* __restrict__ q_scale,\n const float* __restrict__ k_scale,\n const float* __restrict__ v_scale,\n float* __restrict__ out,\n int batch_size,\n int seq_len,\n int num_heads,\n int head_dim,\n float scale\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = batch_size * num_heads * seq_len;\n if (idx >= total) return;\n\n int t = idx % seq_len;\n int h = (idx / seq_len) % num_heads;\n int b = idx / (seq_len * num_heads);\n\n float q_scale_h = q_scale[h];\n float k_scale_h = k_scale[h];\n float v_scale_h = v_scale[h];\n\n float scores[72];\n float max_score = -1e30f;\n for (int s = 0; s <= t; ++s) {\n float dot = 0.0f;\n for (int d = 0; d < head_dim; ++d) {\n float q_val = static_cast<float>(q_q[((b * seq_len + t) * num_heads + h) * head_dim + d]) * q_scale_h;\n float k_val = static_cast<float>(k_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * k_scale_h;\n dot += q_val * k_val;\n }\n scores[s] = dot * scale;\n if (scores[s] > max_score) max_score = scores[s];\n }\n\n float sum_exp = 0.0f;\n for (int s = 0; s <= t; ++s) {\n scores[s] = expf(scores[s] - max_score);\n sum_exp += scores[s];\n }\n float inv_sum = 1.0f / sum_exp;\n\n for (int d = 0; d < head_dim; ++d) {\n float val = 0.0f;\n for (int s = 0; s <= t; ++s) {\n float v_val = static_cast<float>(v_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * v_scale_h;\n val += scores[s] * inv_sum * v_val;\n }\n out[((b * seq_len + t) * num_heads + h) * head_dim + d] = val;\n }\n}\n\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n) {\n auto batch_size = q_q.size(0);\n auto seq_len = q_q.size(1);\n auto num_heads = q_q.size(2);\n auto head_dim = q_q.size(3);\n\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(q_q.device());\n auto out = torch::zeros({batch_size, seq_len, num_heads, head_dim}, options);\n\n float scale = 1.0f / sqrtf((float)head_dim);\n\n int total = batch_size * num_heads * seq_len;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n causal_attention_kernel<<<blocks, threads>>>(\n q_q.data_ptr<int8_t>(),\n k_q.data_ptr<int8_t>(),\n v_q.data_ptr<int8_t>(),\n q_scale.data_ptr<float>(),\n k_scale.data_ptr<float>(),\n v_scale.data_ptr<float>(),\n out.data_ptr<float>(),\n batch_size,\n seq_len,\n num_heads,\n head_dim,\n scale\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"causal_attention\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"causal_attention_cuda\"],\n verbose=False,\n extra_cuda_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, q_q, k_q, v_q, q_scale, k_scale, v_scale):\n device = q_q.device\n q_q = q_q.cuda().contiguous()\n k_q = k_q.cuda().contiguous()\n v_q = v_q.cuda().contiguous()\n q_scale = q_scale.cuda().contiguous()\n k_scale = k_scale.cuda().contiguous()\n v_scale = v_scale.cuda().contiguous()\n out = self.module.causal_attention_cuda(q_q, k_q, v_q, q_scale, k_scale, v_scale)\n return out.to(device)\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: ? trials did not all match.
runtime_error: expected scalar type Char but found Float
turn 5391.3s turn · 10.7s LLM · 380.6s tools · tokens in=9,251 · out=1,087 · total=10,843 · FINAL
reasoning
Still failing. The error must be that one of the tensors passed to the C++ function is float instead of int8.
I was troubleshooting a PyTorch model where the forward computation needed to be replaced with a custom CUDA kernel. The goal was to create a ModelNew class that mirrors the original Model's constructor and forward signature, ensuring numerically equivalent outputs.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void causal_attention_kernel(\n const int8_t* __restrict__ q_q,\n const int8_t* __restrict__ k_q,\n const int8_t* __restrict__ v_q,\n const float* __restrict__ q_scale,\n const float* __restrict__ k_scale,\n const float* __restrict__ v_scale,\n float* __restrict__ out,\n int batch_size,\n int seq_len,\n int num_heads,\n int head_dim,\n float scale\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = batch_size * num_heads * seq_len;\n if (idx >= total) return;\n\n int t = idx % seq_len;\n int h = (idx / seq_len) % num_heads;\n int b = idx / (seq_len * num_heads);\n\n float q_scale_h = q_scale[h];\n float k_scale_h = k_scale[h];\n float v_scale_h = v_scale[h];\n\n float scores[72];\n float max_score = -1e30f;\n for (int s = 0; s <= t; ++s) {\n float dot = 0.0f;\n for (int d = 0; d < head_dim; ++d) {\n float q_val = static_cast<float>(q_q[((b * seq_len + t) * num_heads + h) * head_dim + d]) * q_scale_h;\n float k_val = static_cast<float>(k_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * k_scale_h;\n dot += q_val * k_val;\n }\n scores[s] = dot * scale;\n if (scores[s] > max_score) max_score = scores[s];\n }\n\n float sum_exp = 0.0f;\n for (int s = 0; s <= t; ++s) {\n scores[s] = expf(scores[s] - max_score);\n sum_exp += scores[s];\n }\n float inv_sum = 1.0f / sum_exp;\n\n for (int d = 0; d < head_dim; ++d) {\n float val = 0.0f;\n for (int s = 0; s <= t; ++s) {\n float v_val = static_cast<float>(v_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * v_scale_h;\n val += scores[s] * inv_sum * v_val;\n }\n out[((b * seq_len + t) * num_heads + h) * head_dim + d] = val;\n }\n}\n\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n) {\n auto batch_size = q_q.size(0);\n auto seq_len = q_q.size(1);\n auto num_heads = q_q.size(2);\n auto head_dim = q_q.size(3);\n\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(q_q.device());\n auto out = torch::zeros({batch_size, seq_len, num_heads, head_dim}, options);\n\n float scale = 1.0f / sqrtf((float)head_dim);\n\n int total = batch_size * num_heads * seq_len;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n causal_attention_kernel<<<blocks, threads>>>(\n q_q.data_ptr<int8_t>(),\n k_q.data_ptr<int8_t>(),\n v_q.data_ptr<int8_t>(),\n q_scale.data_ptr<float>(),\n k_scale.data_ptr<float>(),\n v_scale.data_ptr<float>(),\n out.data_ptr<float>(),\n batch_size,\n seq_len,\n num_heads,\n head_dim,\n scale\n );\n\n return out;\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor causal_attention_cuda(\n torch::Tensor q_q,\n torch::Tensor k_q,\n torch::Tensor v_q,\n torch::Tensor q_scale,\n torch::Tensor k_scale,\n torch::Tensor v_scale\n);\n\"\"\"\n\nmodule = load_inline(\n name=\"causal_attention\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"causal_attention_cuda\"],\n verbose=False,\n extra_cuda_cflags=[\"-O3\"]\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.module = module\n\n def forward(self, q_q, k_q, v_q, q_scale, k_scale, v_scale):\n return self.module.causal_attention_cuda(q_q, k_q, v_q, q_scale, k_scale, v_scale)\n"
}
result ← submit_kernel (fail)
failsubmit_kernel
submit_kernel FAILED: correctness check did not pass (? trials).
runtime_error: expected scalar type Char but found Float
final submitted kernel (135 lines, 3,751 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = """
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
__global__ void causal_attention_kernel(
const int8_t* __restrict__ q_q,
const int8_t* __restrict__ k_q,
const int8_t* __restrict__ v_q,
const float* __restrict__ q_scale,
const float* __restrict__ k_scale,
const float* __restrict__ v_scale,
float* __restrict__ out,
int batch_size,
int seq_len,
int num_heads,
int head_dim,
float scale
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = batch_size * num_heads * seq_len;
if (idx >= total) return;
int t = idx % seq_len;
int h = (idx / seq_len) % num_heads;
int b = idx / (seq_len * num_heads);
float q_scale_h = q_scale[h];
float k_scale_h = k_scale[h];
float v_scale_h = v_scale[h];
float scores[72];
float max_score = -1e30f;
for (int s = 0; s <= t; ++s) {
float dot = 0.0f;
for (int d = 0; d < head_dim; ++d) {
float q_val = static_cast<float>(q_q[((b * seq_len + t) * num_heads + h) * head_dim + d]) * q_scale_h;
float k_val = static_cast<float>(k_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * k_scale_h;
dot += q_val * k_val;
}
scores[s] = dot * scale;
if (scores[s] > max_score) max_score = scores[s];
}
float sum_exp = 0.0f;
for (int s = 0; s <= t; ++s) {
scores[s] = expf(scores[s] - max_score);
sum_exp += scores[s];
}
float inv_sum = 1.0f / sum_exp;
for (int d = 0; d < head_dim; ++d) {
float val = 0.0f;
for (int s = 0; s <= t; ++s) {
float v_val = static_cast<float>(v_q[((b * seq_len + s) * num_heads + h) * head_dim + d]) * v_scale_h;
val += scores[s] * inv_sum * v_val;
}
out[((b * seq_len + t) * num_heads + h) * head_dim + d] = val;
}
}
torch::Tensor causal_attention_cuda(
torch::Tensor q_q,
torch::Tensor k_q,
torch::Tensor v_q,
torch::Tensor q_scale,
torch::Tensor k_scale,
torch::Tensor v_scale
) {
auto batch_size = q_q.size(0);
auto seq_len = q_q.size(1);
auto num_heads = q_q.size(2);
auto head_dim = q_q.size(3);
auto options = torch::TensorOptions().dtype(torch::kFloat32).device(q_q.device());
auto out = torch::zeros({batch_size, seq_len, num_heads, head_dim}, options);
float scale = 1.0f / sqrtf((float)head_dim);
int total = batch_size * num_heads * seq_len;
int threads = 256;
int blocks = (total + threads - 1) / threads;
causal_attention_kernel<<<blocks, threads>>>(
q_q.data_ptr<int8_t>(),
k_q.data_ptr<int8_t>(),
v_q.data_ptr<int8_t>(),
q_scale.data_ptr<float>(),
k_scale.data_ptr<float>(),
v_scale.data_ptr<float>(),
out.data_ptr<float>(),
batch_size,
seq_len,
num_heads,
head_dim,
scale
);
return out;
}
"""
cpp_source = """
torch::Tensor causal_attention_cuda(
torch::Tensor q_q,
torch::Tensor k_q,
torch::Tensor v_q,
torch::Tensor q_scale,
torch::Tensor k_scale,
torch::Tensor v_scale
);
"""
module = load_inline(
name="causal_attention",
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=["causal_attention_cuda"],
verbose=False,
extra_cuda_cflags=["-O3"]
)
class ModelNew(nn.Module):
def __init__(self):
super().__init__()
self.module = module
def forward(self, q_q, k_q, v_q, q_scale, k_scale, v_scale):
return self.module.causal_attention_cuda(q_q, k_q, v_q, q_scale, k_scale, v_scale)