can not support 2080ti

[bltcn]

i see this prhttps://github.com/flashinfer-ai/flashinfer/pull/109 is already support sm75. but i use sglang it have error as following:
[gpu_id=0] Prefill batch. #new-seq: 1, #new-token: 6, #cached-token: 0, cache hit rate: 0.00%, #running-req: 0, #queue-req: 0
CUDA Error: no kernel image is available for execution on the device (209) /tmp/build-via-sdist-akw2qk94/flashinfer-0.1.1+cu121torch2.3/include/flashinfer/attention/prefill.cuh: line 2128 at function cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size)
CUDA Error: CUDA Error: no kernel image is available for execution on the device (no kernel image is available for execution on the device209 () 209/tmp/build-via-sdist-akw2qk94/flashinfer-0.1.1+cu121torch2.3/include/flashinfer/attention/prefill.cuh) : line /tmp/build-via-sdist-akw2qk94/flashinfer-0.1.1+cu121torch2.3/include/flashinfer/attention/prefill.cuh2128: line at function 2128cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size) at function
cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size)
Exception in ModelTpServer:
Traceback (most recent call last):
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/tp_worker.py", line 186, in exposed_step
self.forward_step()
File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context
return func(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/tp_worker.py", line 202, in forward_step
self.forward_prefill_batch(new_batch)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/tp_worker.py", line 443, in forward_prefill_batch
output = self.model_runner.forward(batch, ForwardMode.EXTEND)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/model_runner.py", line 336, in forward
return self.forward_extend(batch)
File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context
return func(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/model_runner.py", line 304, in forward_extend
return self.model.forward(
File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context
return func(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/models/qwen2.py", line 272, in forward
hidden_states = self.model(input_ids, positions, input_metadata, input_embeds)
File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
return self._call_impl(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl
return forward_call(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/models/qwen2.py", line 240, in forward
hidden_states, residual = layer(
File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
return self._call_impl(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl
return forward_call(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/models/qwen2.py", line 192, in forward
hidden_states = self.self_attn(
File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
return self._call_impl(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl
return forward_call(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/models/qwen2.py", line 141, in forward
attn_output = self.attn(q, k, v, input_metadata)
File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
return self._call_impl(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl
return forward_call(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/layers/radix_attention.py", line 149, in forward
return self.extend_forward(q, k, v, input_metadata)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/layers/radix_attention.py", line 91, in extend_forward_flashinfer
o = input_metadata.flashinfer_prefill_wrapper_paged.forward(
File "/usr/local/lib/python3.10/dist-packages/flashinfer/prefill.py", line 875, in forward
return self._wrapper.forward(
RuntimeError: BatchPrefillWithPagedKVCache failed with error code no kernel image is available for execution on the device

[yzh119]

2080ti has cuda arch of sm75 which we didn't set it as a compilation target in our pre-built wheels.
You can try compile the kernels from source code

cd python
pip install -e . -v

I think most of the kernels should work on sm75 but I haven't fully tested them.

[bltcn]

thanks. i'll try it

[bltcn]

@yzh119 other question, which TORCH_CUDA_ARCH_LIST is I can set?

[bltcn]

I compile the source code and get this:
[gpu_id=0] Prefill batch. #new-seq: 1, #new-token: 6, #cached-token: 0, cache hit rate: 0.00%, #running-req: 0, #queue-req: 0
CUDA Error: CUDA Error: too many resources requested for launch (701) /sgl-workspace/flashinfer/python/include/flashinfer/sampling.cuh: line 547 at function cudaLaunchKernel((void*)kernel, nblks, nthrs, args, smem_size, stream)
CUDA Error: too many resources requested for launch (701) /sgl-workspace/flashinfer/python/include/flashinfer/sampling.cuh: line 547 at function too many resources requested for launchcudaLaunchKernel((void*)kernel, nblks, nthrs, args, smem_size, stream) (
701) /sgl-workspace/flashinfer/python/include/flashinfer/sampling.cuh: line 547 at function cudaLaunchKernel((void*)kernel, nblks, nthrs, args, smem_size, stream)
CUDA Error: too many resources requested for launch (701) /sgl-workspace/flashinfer/python/include/flashinfer/sampling.cuh: line 547 at function cudaLaunchKernel((void*)kernel, nblks, nthrs, args, smem_size, stream)
Exception in ModelTpServer:
Traceback (most recent call last):
File "/sgl-workspace/sglang/python/sglang/srt/managers/controller/tp_worker.py", line 186, in exposed_step
self.forward_step()
File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context
return func(*args, **kwargs)
File "/sgl-workspace/sglang/python/sglang/srt/managers/controller/tp_worker.py", line 202, in forward_step
self.forward_prefill_batch(new_batch)
File "/sgl-workspace/sglang/python/sglang/srt/managers/controller/tp_worker.py", line 444, in forward_prefill_batch
next_token_ids = batch.sample(output.next_token_logits)
File "/sgl-workspace/sglang/python/sglang/srt/managers/controller/infer_batch.py", line 692, in sample
batch_next_token_ids, success = top_k_top_p_sampling_from_probs(
File "/sgl-workspace/flashinfer/python/flashinfer/sampling.py", line 266, in top_k_top_p_sampling_from_probs
return _kernels.top_k_top_p_sampling_from_probs(
RuntimeError: TopKTopPSamplingFromProbs failed with error code too many resources requested for launch

[yzh119]

I think it's because sm75 have small shared memory size.

Would you mind trying setting BLOCK_THREADS to some smaller values such as 256 or 512, and compile again?

flashinfer/include/flashinfer/sampling.cuh

Line 534 in 73a764b

constexpr uint32_t BLOCK_THREADS = 1024;

[bltcn]

@yzh119 ok,thanks. I try it.

[bltcn]

@yzh119 It dosen't work.

[gpu_id=0] Prefill batch. #new-seq: 1, #new-token: 6, #cached-token: 0, cache hit rate: 0.00%, #running-req: 0, #queue-req: 0
Exception in ModelTpServer:
Traceback (most recent call last):
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/tp_worker.py", line 186, in exposed_step
self.forward_step()
File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context
return func(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/tp_worker.py", line 216, in forward_step
self.forward_decode_batch(self.running_batch)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/tp_worker.py", line 567, in forward_decode_batch
output = self.model_runner.forward(batch, ForwardMode.DECODE)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/model_runner.py", line 334, in forward
return self.forward_decode(batch)
File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context
return func(*args, **kwargs)
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/model_runner.py", line 276, in forward_decode
input_metadata = InputMetadata.create(
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/infer_batch.py", line 770, in create
init_flashinfer_args(
File "/usr/local/lib/python3.10/dist-packages/sglang/srt/managers/controller/infer_batch.py", line 884, in init_flashinfer_args
flashinfer_decode_wrapper.begin_forward(
File "/sgl-workspace/flashinfer/python/flashinfer/decode.py", line 517, in begin_forward
self._wrapper.begin_forward(
RuntimeError: BatchDecodeWithPagedKVCachePyTorchWrapper::BeginForward(at::Tensor, at::Tensor, at::Tensor, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, float, at::Tensor, at::Tensor)::<lambda()> failed to dispatch fp8 data type Half

[yzh119]

@bltcn I think it's because you turned off enable_fp8:

flashinfer/python/setup.py

Line 38 in 2496f5b

enable_fp8 = True

Set this value to True and then this error message should have gone.

[bltcn]

@bltcn I think it's because you turned off enable_fp8:

flashinfer/python/setup.py

Line 38 in 2496f5b

enable_fp8 = True

Set this value to True and then this error message should have gone.

can't complie.

[80/400] /usr/local/cuda/bin/nvcc --generate-dependencies-with-compile --dependency-output /sgl-workspace/flashinfer/python/build/temp.linux-x86_64-3.10/csrc/generated/batch_paged_prefill_head_256_logitshook_0_posenc_1_fp16qkred_0_mask_2_dtypein_f16_dtypeout_f16_idtype_i32.o.d -I/sgl-workspace/flashinfer/python/include -I/sgl-workspace/flashinfer/python/3rdparty/cutlass/include -I/usr/local/lib/python3.10/dist-packages/torch/include -I/usr/local/lib/python3.10/dist-packages/torch/include/torch/csrc/api/include -I/usr/local/lib/python3.10/dist-packages/torch/include/TH -I/usr/local/lib/python3.10/dist-packages/torch/include/THC -I/usr/local/cuda/include -I/usr/include/python3.10 -c -c /sgl-workspace/flashinfer/python/csrc/generated/batch_paged_prefill_head_256_logitshook_0_posenc_1_fp16qkred_0_mask_2_dtypein_f16_dtypeout_f16_idtype_i32.cu -o /sgl-workspace/flashinfer/python/build/temp.linux-x86_64-3.10/csrc/generated/batch_paged_prefill_head_256_logitshook_0_posenc_1_fp16qkred_0_mask_2_dtypein_f16_dtypeout_f16_idtype_i32.o --expt-relaxed-constexpr --compiler-options ''"'"'-fPIC'"'"'' -O3 -std=c++17 --threads 1 -Xfatbin -compress-all -DTORCH_API_INCLUDE_EXTENSION_H '-DPYBIND11_COMPILER_TYPE="_gcc"' '-DPYBIND11_STDLIB="_libstdcpp"' '-DPYBIND11_BUILD_ABI="_cxxabi1011"' -DTORCH_EXTENSION_NAME=_kernels -D_GLIBCXX_USE_CXX11_ABI=0 -gencode=arch=compute_75,code=sm_75
ninja: build stopped: subcommand failed.
Traceback (most recent call last):
File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 2107, in _run_ninja_build
subprocess.run(
File "/usr/lib/python3.10/subprocess.py", line 526, in run
raise CalledProcessError(retcode, process.args,
subprocess.CalledProcessError: Command '['ninja', '-v', '-j', '80']' returned non-zero exit status 1.

The above exception was the direct cause of the following exception:

Traceback (most recent call last):
  File "<string>", line 2, in <module>
  File "<pip-setuptools-caller>", line 34, in <module>
  File "/sgl-workspace/flashinfer/python/setup.py", line 349, in <module>
    setuptools.setup(
  File "/usr/lib/python3/dist-packages/setuptools/__init__.py", line 153, in setup
    return distutils.core.setup(**attrs)
  File "/usr/lib/python3.10/distutils/core.py", line 148, in setup
    dist.run_commands()
  File "/usr/lib/python3.10/distutils/dist.py", line 966, in run_commands
    self.run_command(cmd)
  File "/usr/lib/python3.10/distutils/dist.py", line 985, in run_command
    cmd_obj.run()
  File "/usr/lib/python3/dist-packages/setuptools/command/develop.py", line 34, in run
    self.install_for_development()
  File "/usr/lib/python3/dist-packages/setuptools/command/develop.py", line 114, in install_for_development
    self.run_command('build_ext')
  File "/usr/lib/python3.10/distutils/cmd.py", line 313, in run_command
    self.distribution.run_command(command)
  File "/usr/lib/python3.10/distutils/dist.py", line 985, in run_command
    cmd_obj.run()
  File "/usr/lib/python3/dist-packages/setuptools/command/build_ext.py", line 79, in run
    _build_ext.run(self)
  File "/usr/lib/python3.10/distutils/command/build_ext.py", line 340, in run
    self.build_extensions()
  File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 870, in build_extensions
    build_ext.build_extensions(self)
  File "/usr/lib/python3.10/distutils/command/build_ext.py", line 449, in build_extensions
    self._build_extensions_serial()
  File "/usr/lib/python3.10/distutils/command/build_ext.py", line 474, in _build_extensions_serial
    self.build_extension(ext)
  File "/usr/lib/python3/dist-packages/setuptools/command/build_ext.py", line 202, in build_extension
    _build_ext.build_extension(self, ext)
  File "/usr/lib/python3.10/distutils/command/build_ext.py", line 529, in build_extension
    objects = self.compiler.compile(sources,
  File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 683, in unix_wrap_ninja_compile
    _write_ninja_file_and_compile_objects(
  File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 1783, in _write_ninja_file_and_compile_objects
    _run_ninja_build(
  File "/usr/local/lib/python3.10/dist-packages/torch/utils/cpp_extension.py", line 2123, in _run_ninja_build
    raise RuntimeError(message) from e
RuntimeError: Error compiling objects for extension
error: subprocess-exited-with-error

× python setup.py develop did not run successfully.
│ exit code: 1
╰─> See above for output.

note: This error originates from a subprocess, and is likely not a problem with pip.
full command: /usr/bin/python3 -c '
exec(compile('"'"''"'"''"'"'
# This is <pip-setuptools-caller> -- a caller that pip uses to run setup.py
#
# - It imports setuptools before invoking setup.py, to enable projects that directly
#   import from `distutils.core` to work with newer packaging standards.
# - It provides a clear error message when setuptools is not installed.
# - It sets `sys.argv[0]` to the underlying `setup.py`, when invoking `setup.py` so
#   setuptools doesn'"'"'t think the script is `-c`. This avoids the following warning:
#     manifest_maker: standard file '"'"'-c'"'"' not found".
# - It generates a shim setup.py, for handling setup.cfg-only projects.
import os, sys, tokenize

try:
    import setuptools
except ImportError as error:
    print(
        "ERROR: Can not execute `setup.py` since setuptools is not available in "
        "the build environment.",
        file=sys.stderr,
    )
    sys.exit(1)

__file__ = %r
sys.argv[0] = __file__

if os.path.exists(__file__):
    filename = __file__
    with tokenize.open(__file__) as f:
        setup_py_code = f.read()
else:
    filename = "<auto-generated setuptools caller>"
    setup_py_code = "from setuptools import setup; setup()"

exec(compile(setup_py_code, filename, "exec"))
'"'"''"'"''"'"' % ('"'"'/sgl-workspace/flashinfer/python/setup.py'"'"',), "<pip-setuptools-caller>", "exec"))' develop --no-deps
cwd: /sgl-workspace/flashinfer/python/

Rolling back uninstall of flashinfer
Moving to /usr/local/lib/python3.10/dist-packages/flashinfer-0.1.1+cu121torch2.3.dist-info/
from /usr/local/lib/python3.10/dist-packages/~lashinfer-0.1.1+cu121torch2.3.dist-info
Moving to /usr/local/lib/python3.10/dist-packages/flashinfer/
from /usr/local/lib/python3.10/dist-packages/~lashinfer
error: subprocess-exited-with-error

× python setup.py develop did not run successfully.
│ exit code: 1
╰─> See above for output.

note: This error originates from a subprocess, and is likely not a problem with pip.

[notice] A new release of pip is available: 24.1.2 -> 24.2
[notice] To update, run: python3 -m pip install --upgrade pip

[yzh119]

Hi @bltcn , the error message is not complete and I cannot identify the potential issue.
Would you mind posting more error messages that includes the nvcc/g++ error (seems they are truncated)?

I don't have a sm75 machine so it's hard for me to debug without full error log.

[bltcn]

Hi @bltcn , the error message is not complete and I cannot identify the potential issue. Would you mind posting more error messages that includes the nvcc/g++ error (seems they are truncated)?

I don't have a sm75 machine so it's hard for me to debug without full error log.

i can proide one machine with 8*2080ti(22g), if you need. Please contact me.

[yzh119]

Please send me an email(zhye@cs.washington.edu) about the server access, and I'll figure out how to fix it.

[bltcn]

Please send me an email(zhye@cs.washington.edu) about the server access, and I'll figure out how to fix it.

I had sent the server info to you yesterday by email(bltcn@163.com). If you have some problem, please tell me.