XYZboom opened a new issue, #20106:
URL: https://github.com/apache/tvm/issues/20106
### Expected behavior
`relax.op.cumprod` is a valid Relax operator. `relax.build(mod,
target="cuda")` should compile successfully, and the compiled module should
execute and return correct cumulative-product results for any valid input shape.
### Actual behavior
The CUDA `cumprod` kernel launch throws `CUDALaunch Error:
CUDA_ERROR_INVALID_VALUE`:
```
tvm.error.InternalError: CUDALaunch Error: CUDA_ERROR_INVALID_VALUE
grid=(1,100800,1), block=(1024,1,1)
// func_name=cumprod_kernel
```
The crash happens at **execution time** (kernel launch), not at build time.
The kernel launch config `grid=(1, rows, 1)` puts `gridDim.y = rows`, which
exceeds the CUDA hardware limit of 65535, so the launch fails. The error is
deterministic — the same input shape always triggers it.
### Environment
- **OS**: Linux (x86_64, conda environment)
- **GPU**: NVIDIA GeForce RTX 3080 Ti (12GB VRAM)
- **CUDA driver**: 580.76.05
- **TVM version**: 0.25.0.post1 (pip install)
- **Target**: `cuda`
- **Python**: 3.12.13
### Steps to reproduce
```python
import tvm
from tvm import relax
import numpy as np
bb = relax.BlockBuilder()
v = relax.Var("v", relax.TensorStructInfo(
shape=relax.ShapeExpr([65536, 1]), dtype="float32"))
with bb.function("main", [v]):
out = bb.emit(relax.op.cumprod(v, axis=1))
bb.emit_func_output(out)
mod = bb.get()
ex = relax.build(mod, target="cuda")
vm = relax.VirtualMachine(ex, tvm.cuda())
np_in = np.random.uniform(0.0, 1.0, size=(65536, 1)).astype(np.float32)
t_in = tvm.runtime.tensor(np_in, device=tvm.cuda())
result = vm["main"](t_in) # crashes here
```
error log:
```txt
Traceback (most recent call last):
File
"/root/autodl-tmp/data/maybeBug/bug_014_TVM_Relax_(daemon)_seed20261649_1_20260808_095843/reduced.py",
line 24, in <module>
result = vm["main"](t_in)
^^^^^^^^^^^^^^^^
File "python/tvm_ffi/cython/function.pxi", line 968, in
tvm_ffi.core.Function.__call__
File "/__w/tvm/tvm/src/backend/cuda/runtime/cuda_module.cc", line 305, in
void tvm::runtime::CUDAWrappedFunc::operator()(tvm::ffi::PackedArgs,
tvm::ffi::Any*, void**) const
tvm.error.InternalError: CUDALaunch Error: CUDA_ERROR_INVALID_VALUE
grid=(1,65536,1), block=(1024,1,1)
// func_name=cumprod_kernel
// CUDA Source
// -----------
#ifdef __CUDACC_RTC__
#include <cuda/std/cstdint>
using cuda::std::uint8_t;
using cuda::std::uint16_t;
using cuda::std::uint32_t;
using cuda::std::uint64_t;
using cuda::std::int8_t;
using cuda::std::int16_t;
using cuda::std::int32_t;
using cuda::std::int64_t;
#include <cuda/std/type_traits>
namespace std {
using cuda::std::is_same;
using cuda::std::is_same_v;
using cuda::std::is_integral;
using cuda::std::is_signed;
using cuda::std::is_unsigned;
using cuda::std::is_floating_point;
using cuda::std::enable_if;
using cuda::std::conditional;
}
// NVRTC uses asm/volatile instead of __asm__/__volatile__ (gcc extension).
#ifndef __asm__
#define __asm__ asm
#endif
#ifndef __volatile__
#define __volatile__ volatile
#endif
#else
#include <cstdint>
#include <type_traits>
#include <cuda.h>
#endif
#if (((__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ >= 4)) || \
(__CUDACC_VER_MAJOR__ > 11))
#define TVM_ENABLE_L2_PREFETCH 1
#else
#define TVM_ENABLE_L2_PREFETCH 0
#endif
#ifdef _WIN32
using uint = unsigned int;
using uchar = unsigned char;
using ushort = unsigned short;
using int64_t = long long;
using uint64_t = unsigned long long;
#else
#define uint unsigned int
#define uchar unsigned char
#define ushort unsigned short
#endif
extern "C" __global__ void __launch_bounds__(1024) cumprod_kernel(float*
__restrict__ data_buf, float* __restrict__ output_buf);
extern "C" __global__ void cumprod_kernel_1(float* __restrict__ output_buf);
extern "C" __global__ void __launch_bounds__(1024) cumprod_kernel_2(float*
__restrict__ T_multiply, float* __restrict__ data_buf, float* __restrict__
output_buf);
extern "C" __global__ void __launch_bounds__(1024) cumprod_kernel(float*
__restrict__ data_buf, float* __restrict__ output_buf) {
if (((int)threadIdx.x) < 1) {
output_buf[((int)blockIdx.y)] = data_buf[((int)blockIdx.y)];
}
}
extern "C" __global__ void cumprod_kernel_1(float* __restrict__ output_buf) {
output_buf[((int)blockIdx.x)] = 0x1p+0f/*1.000000e+00*/;
}
extern "C" __global__ void __launch_bounds__(1024) cumprod_kernel_2(float*
__restrict__ T_multiply, float* __restrict__ data_buf, float* __restrict__
output_buf) {
int cse_v1 = ((((int)blockIdx.x) * 1024) + ((int)threadIdx.x));
T_multiply[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] =
(data_buf[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))] *
output_buf[((((int)blockIdx.x) * 1024) + ((int)threadIdx.x))]);
}
```
### Triage
* backend:cuda
* cumprod
* needs-triage
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
