deepseek-ai/DeepGEMM
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepGEMM.git synced 2026-11-05 00:47:26 +00:00
Code Issues 162 Projects Packages Wiki Activity

[GH-ISSUE #106] Bug: AssertionError in gemm_fp8_fp8_bf16_nt due to new test case in test_core.py #60

New issue
Closed
opened 2026-10-26 21:11:30 +00:00 by bsakaguchi · 8 comments
bsakaguchi commented 2026-10-26 21:11:30 +00:00
Owner
Copy link

Originally created by @yhyang201 on GitHub (May 23, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/106

Tests are failing with an AssertionError in the gemm_fp8_fp8_bf16_nt function. This issue was introduced after merging PR #95 (commit 04278f6dee).

The failure is caused by a new test case (k=576, n=7168) added at line 167 in tests/test_core.py.

        for k, n in [(576, 7168), (7168, 2112), (1536, 24576), (512, 32768), (16384, 7168), (7168, 4096), (2048, 7168)]:

Testing GEMM:
Traceback (most recent call last):
  File "/home/ubuntu/DeepGEMM/tests/test_core.py", line 310, in <module>
    test_gemm()
  File "/home/ubuntu/DeepGEMM/tests/test_core.py", line 169, in test_gemm
    deep_gemm.gemm_fp8_fp8_bf16_nt(x_fp8, y_fp8, out)
  File "/home/ubuntu/.python/sglang/lib/python3.10/site-packages/deep_gemm/jit_kernels/gemm.py", line 167, in gemm_fp8_fp8_bf16_nt
    assert n % 64 == 0 and k % 128 == 0
AssertionError
Originally created by @yhyang201 on GitHub (May 23, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/106 Tests are failing with an `AssertionError` in the `gemm_fp8_fp8_bf16_nt` function. This issue was introduced after merging PR #95 (commit 04278f6dee2a7b9da5d881fa40b24ad27dfac0f6). The failure is caused by a new test case `(k=576, n=7168)` added at line 167 in `tests/test_core.py`. ``` for k, n in [(576, 7168), (7168, 2112), (1536, 24576), (512, 32768), (16384, 7168), (7168, 4096), (2048, 7168)]: ``` ``` Testing GEMM: Traceback (most recent call last): File "/home/ubuntu/DeepGEMM/tests/test_core.py", line 310, in <module> test_gemm() File "/home/ubuntu/DeepGEMM/tests/test_core.py", line 169, in test_gemm deep_gemm.gemm_fp8_fp8_bf16_nt(x_fp8, y_fp8, out) File "/home/ubuntu/.python/sglang/lib/python3.10/site-packages/deep_gemm/jit_kernels/gemm.py", line 167, in gemm_fp8_fp8_bf16_nt assert n % 64 == 0 and k % 128 == 0 AssertionError ```
bsakaguchi commented 2026-10-26 21:11:31 +00:00
Author
Owner
Copy link

@yhyang201 commented on GitHub (May 23, 2025):

The test_jit.py file also has an issue. Here's the generated code:

#ifdef __CUDACC_RTC__
#include <deep_gemm/nvrtc_std.cuh>
#else
#include <cuda.h>
#endif

#include <cuda_fp8.h>
#include <cuda_bf16.h>

template <typename T>
__global__ void vector_add(T* a, T* b, T* c, uint32_t n) {
    uint32_t i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

static void __instantiate_kernel() {
    auto ptr = reinterpret_cast<void*>(&vector_add<float>);
}

Compiler: NVCC, version: (12, 9)
Building ...
Traceback (most recent call last):
File "/home/ubuntu/DeepGEMM/tests/test\_jit.py", line 89, in <module>
func = compiler\_cls.build('test\_func', code, VectorAddRuntime, kwargs)
TypeError: Compiler.build() takes 4 positional arguments but 5 were given
<!-- gh-comment-id:2905115919 --> @yhyang201 commented on GitHub (May 23, 2025): The `test_jit.py` file also has an issue. Here's the generated code: ``` #ifdef __CUDACC_RTC__ #include <deep_gemm/nvrtc_std.cuh> #else #include <cuda.h> #endif #include <cuda_fp8.h> #include <cuda_bf16.h> template <typename T> __global__ void vector_add(T* a, T* b, T* c, uint32_t n) { uint32_t i = blockDim.x * blockIdx.x + threadIdx.x; if (i < n) { c[i] = a[i] + b[i]; } } static void __instantiate_kernel() { auto ptr = reinterpret_cast<void*>(&vector_add<float>); } Compiler: NVCC, version: (12, 9) Building ... Traceback (most recent call last): File "/home/ubuntu/DeepGEMM/tests/test\_jit.py", line 89, in <module> func = compiler\_cls.build('test\_func', code, VectorAddRuntime, kwargs) TypeError: Compiler.build() takes 4 positional arguments but 5 were given ```
bsakaguchi commented 2026-10-26 21:11:32 +00:00
Author
Owner
Copy link

@yhyang201 commented on GitHub (May 23, 2025):

At the same time, I encountered the following issue with commit 04278f6dee:

Testing weight gradient GEMM:
Traceback (most recent call last):
  File "/home/ubuntu/DeepGEMM/tests/test_core.py", line 314, in <module>
    test_wgrad_gemm()
  File "/home/ubuntu/DeepGEMM/tests/test_core.py", line 252, in test_wgrad_gemm
    deep_gemm.wgrad_gemm_fp8_fp8_fp32_nt(x_fp8, y_fp8, out)
AttributeError: module 'deep_gemm' has no attribute 'wgrad_gemm_fp8_fp8_fp32_nt'
<!-- gh-comment-id:2905172725 --> @yhyang201 commented on GitHub (May 23, 2025): At the same time, I encountered the following issue with commit 04278f6dee2a7b9da5d881fa40b24ad27dfac0f6: ``` Testing weight gradient GEMM: Traceback (most recent call last): File "/home/ubuntu/DeepGEMM/tests/test_core.py", line 314, in <module> test_wgrad_gemm() File "/home/ubuntu/DeepGEMM/tests/test_core.py", line 252, in test_wgrad_gemm deep_gemm.wgrad_gemm_fp8_fp8_fp32_nt(x_fp8, y_fp8, out) AttributeError: module 'deep_gemm' has no attribute 'wgrad_gemm_fp8_fp8_fp32_nt' ```
bsakaguchi commented 2026-10-26 21:11:32 +00:00
Author
Owner
Copy link

@soundOfDestiny commented on GitHub (May 24, 2025):

Thanks!

assert n % 64 == 0 and k % 128 == 0

It's fixed in latest main.

<!-- gh-comment-id:2906329219 --> @soundOfDestiny commented on GitHub (May 24, 2025): Thanks! > assert n % 64 == 0 and k % 128 == 0 It's fixed in latest main.
bsakaguchi commented 2026-10-26 21:11:32 +00:00
Author
Owner
Copy link

@shanguanma commented on GitHub (May 28, 2025):

At the same time, I encountered the following issue when I run tests/test_jit.py


#ifdef __CUDACC_RTC__
#include <deep_gemm/nvrtc_std.cuh>
#else
#include <cuda.h>
#endif

#include <cuda_fp8.h>
#include <cuda_bf16.h>

template <typename T>
__global__ void vector_add(T* a, T* b, T* c, uint32_t n) {
    uint32_t i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

static void __instantiate_kernel() {
    auto ptr = reinterpret_cast<void*>(&vector_add<float>);
}


Compiler: NVCC, version: (12, 4)
Building ...
Compiling JIT runtime kernel.test_func.085a74ce58c7 with command ['/usr/local/cuda-12.4.1/bin/nvcc', '/root/.deep_gemm/cache/kernel.test_func.085a74ce58c7/kernel.cu', '-o', '/root/.deep_gemm/tmp/nvcc.tmp.f598576b-71c4-4af6-95b0-82e3258513da.93e1b2f65ac2.cubin', '-std=c++20', '--ptxas-options=--register-usage-level=10', '--diag-suppress=39,161,174,177,186,940', '-I/maduo/miniconda3/envs/deepgemm/lib/python3.11/site-packages/deep_gemm-1.0.0+8dfa329-py3.11.egg/deep_gemm/jit/../include', '-gencode=arch=compute_90a,code=sm_90a', '-cubin', '-O3', '--expt-relaxed-constexpr', '--expt-extended-lambda', '--compiler-options=-fPIC,-O3,-fconcepts,-Wno-deprecated-declarations,-Wno-abi']
Compilation of JIT runtime kernel.test_func.085a74ce58c7 took 0.74 seconds.
Put kernel kernel.test_func.085a74ce58c7 with {'T': 'float'} into runtime cache
Loading JIT runtime /root/.deep_gemm/cache/kernel.test_func.085a74ce58c7 took 3.89 ms.
ret: 300
Traceback (most recent call last):
  File "/maduo/codebase/DeepGEMM/tests/test_jit.py", line 97, in <module>
    assert ret == cbd.CUresult.CUDA_SUCCESS, ret
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
AssertionError: 300

Do you have any suggestion?

<!-- gh-comment-id:2915508072 --> @shanguanma commented on GitHub (May 28, 2025): At the same time, I encountered the following issue when I run tests/test_jit.py ``` #ifdef __CUDACC_RTC__ #include <deep_gemm/nvrtc_std.cuh> #else #include <cuda.h> #endif #include <cuda_fp8.h> #include <cuda_bf16.h> template <typename T> __global__ void vector_add(T* a, T* b, T* c, uint32_t n) { uint32_t i = blockDim.x * blockIdx.x + threadIdx.x; if (i < n) { c[i] = a[i] + b[i]; } } static void __instantiate_kernel() { auto ptr = reinterpret_cast<void*>(&vector_add<float>); } Compiler: NVCC, version: (12, 4) Building ... Compiling JIT runtime kernel.test_func.085a74ce58c7 with command ['/usr/local/cuda-12.4.1/bin/nvcc', '/root/.deep_gemm/cache/kernel.test_func.085a74ce58c7/kernel.cu', '-o', '/root/.deep_gemm/tmp/nvcc.tmp.f598576b-71c4-4af6-95b0-82e3258513da.93e1b2f65ac2.cubin', '-std=c++20', '--ptxas-options=--register-usage-level=10', '--diag-suppress=39,161,174,177,186,940', '-I/maduo/miniconda3/envs/deepgemm/lib/python3.11/site-packages/deep_gemm-1.0.0+8dfa329-py3.11.egg/deep_gemm/jit/../include', '-gencode=arch=compute_90a,code=sm_90a', '-cubin', '-O3', '--expt-relaxed-constexpr', '--expt-extended-lambda', '--compiler-options=-fPIC,-O3,-fconcepts,-Wno-deprecated-declarations,-Wno-abi'] Compilation of JIT runtime kernel.test_func.085a74ce58c7 took 0.74 seconds. Put kernel kernel.test_func.085a74ce58c7 with {'T': 'float'} into runtime cache Loading JIT runtime /root/.deep_gemm/cache/kernel.test_func.085a74ce58c7 took 3.89 ms. ret: 300 Traceback (most recent call last): File "/maduo/codebase/DeepGEMM/tests/test_jit.py", line 97, in <module> assert ret == cbd.CUresult.CUDA_SUCCESS, ret ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ AssertionError: 300 ``` Do you have any suggestion?
bsakaguchi commented 2026-10-26 21:11:33 +00:00
Author
Owner
Copy link

@LyricZhao commented on GitHub (May 29, 2025):

@shanguanma You can refer to here: https://nvidia.github.io/cuda-python/cuda-bindings/latest/module/driver.html#cuda.bindings.driver.CUresult.CUDA_ERROR_INVALID_SOURCE. Do you use devices that support sm_90a?

<!-- gh-comment-id:2918400333 --> @LyricZhao commented on GitHub (May 29, 2025): @shanguanma You can refer to here: https://nvidia.github.io/cuda-python/cuda-bindings/latest/module/driver.html#cuda.bindings.driver.CUresult.CUDA_ERROR_INVALID_SOURCE. Do you use devices that support `sm_90a`?
bsakaguchi commented 2026-10-26 21:11:33 +00:00
Author
Owner
Copy link

@shanguanma commented on GitHub (May 29, 2025):

@shanguanma You can refer to here: https://nvidia.github.io/cuda-python/cuda-bindings/latest/module/driver.html#cuda.bindings.driver.CUresult.CUDA_ERROR_INVALID_SOURCE. Do you use devices that support sm_90a?
Thanks for your reply, In my mechine, it contains multiple cuda version(i.e. cuda11.8 and cuda12.4.1), the gpu type is A100-SXM4-80GB, how to specifiy cuda version to compile this DeepGEMM?

<!-- gh-comment-id:2918792207 --> @shanguanma commented on GitHub (May 29, 2025): > [@shanguanma](https://github.com/shanguanma) You can refer to here: https://nvidia.github.io/cuda-python/cuda-bindings/latest/module/driver.html#cuda.bindings.driver.CUresult.CUDA_ERROR_INVALID_SOURCE. Do you use devices that support `sm_90a`? Thanks for your reply, In my mechine, it contains multiple cuda version(i.e. cuda11.8 and cuda12.4.1), the gpu type is A100-SXM4-80GB, how to specifiy cuda version to compile this DeepGEMM?
bsakaguchi commented 2026-10-26 21:11:33 +00:00
Author
Owner
Copy link

@soundOfDestiny commented on GitHub (May 29, 2025):

@shanguanma You can refer to here: https://nvidia.github.io/cuda-python/cuda-bindings/latest/module/driver.html#cuda.bindings.driver.CUresult.CUDA_ERROR_INVALID_SOURCE. Do you use devices that support sm_90a?
Thanks for your reply, In my mechine, it contains multiple cuda version(i.e. cuda11.8 and cuda12.4.1), the gpu type is A100-SXM4-80GB, how to specifiy cuda version to compile this DeepGEMM?

Sadly we need hopper to run DeepGEMM. Ampere is not supported :-(

<!-- gh-comment-id:2918796189 --> @soundOfDestiny commented on GitHub (May 29, 2025): > > [@shanguanma](https://github.com/shanguanma) You can refer to here: https://nvidia.github.io/cuda-python/cuda-bindings/latest/module/driver.html#cuda.bindings.driver.CUresult.CUDA_ERROR_INVALID_SOURCE. Do you use devices that support `sm_90a`? > > Thanks for your reply, In my mechine, it contains multiple cuda version(i.e. cuda11.8 and cuda12.4.1), the gpu type is A100-SXM4-80GB, how to specifiy cuda version to compile this DeepGEMM? Sadly we need hopper to run DeepGEMM. Ampere is not supported :-(
bsakaguchi commented 2026-10-26 21:11:33 +00:00
Author
Owner
Copy link

@shanguanma commented on GitHub (May 29, 2025):

The news is so sad for me.

<!-- gh-comment-id:2918818183 --> @shanguanma commented on GitHub (May 29, 2025): The news is so sad for me.
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
public-sync-20261014
0930-release
nv_dev
fix-mega-moe-fence
public-release-2609
public-release-2607
public-release-2606
mega-update
public-release-260416
ko3n1g/ci/build-wheels
nv_dev_f8e8fb5
nv_dev_950d31a
nv_dev_67fc648
nv_dev_c491439
nv_dev_4ff3f54
v2.1.1.post3
v2.1.1.post2
v2.1.1.post1
v2.1.1
v2.1.0
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
deepseek-ai/DeepGEMM#60
No description provided.