issue/1031 fix T2-1-1

Author: PanZezhong1725

include/infiniop/ops/dequantize_gptq.h

@@ -5,26 +5,26 @@
 
 typedef struct InfiniopDescriptor *infiniopDequantizeGPTQDescriptor_t;
 
-__C __export infiniStatus_t infiniopCreateDequantizeGPTQDescriptor(infiniopHandle_t handle,
-                                                                  infiniopDequantizeGPTQDescriptor_t *desc_ptr,
-                                                                  infiniopTensorDescriptor_t out_desc,
-                                                                  infiniopTensorDescriptor_t qweight_desc,
-                                                                  infiniopTensorDescriptor_t scales_desc,
-                                                                  infiniopTensorDescriptor_t zeros_desc,
-                                                                  infiniopTensorDescriptor_t g_idx_desc); // add g_idx
+__INFINI_C __export infiniStatus_t infiniopCreateDequantizeGPTQDescriptor(infiniopHandle_t handle,
+                                                                          infiniopDequantizeGPTQDescriptor_t *desc_ptr,
+                                                                          infiniopTensorDescriptor_t out_desc,
+                                                                          infiniopTensorDescriptor_t qweight_desc,
+                                                                          infiniopTensorDescriptor_t scales_desc,
+                                                                          infiniopTensorDescriptor_t zeros_desc,
+                                                                          infiniopTensorDescriptor_t g_idx_desc); // add g_idx
 
-__C __export infiniStatus_t infiniopGetDequantizeGPTQWorkspaceSize(infiniopDequantizeGPTQDescriptor_t desc, size_t *size);
+__INFINI_C __export infiniStatus_t infiniopGetDequantizeGPTQWorkspaceSize(infiniopDequantizeGPTQDescriptor_t desc, size_t *size);
 
-__C __export infiniStatus_t infiniopDequantizeGPTQ(infiniopDequantizeGPTQDescriptor_t desc,
-                                                  void *workspace,
-                                                  size_t workspace_size,
-                                                  void *out,
-                                                  const void *qweight,
-                                                  const void *scales,
-                                                  const void *zeros,
-                                                  const void *g_idx,  // add g_idx
-                                                  void *stream);
+__INFINI_C __export infiniStatus_t infiniopDequantizeGPTQ(infiniopDequantizeGPTQDescriptor_t desc,
+                                                          void *workspace,
+                                                          size_t workspace_size,
+                                                          void *out,
+                                                          const void *qweight,
+                                                          const void *scales,
+                                                          const void *zeros,
+                                                          const void *g_idx, // add g_idx
+                                                          void *stream);
 
-__C __export infiniStatus_t infiniopDestroyDequantizeGPTQDescriptor(infiniopDequantizeGPTQDescriptor_t desc);
+__INFINI_C __export infiniStatus_t infiniopDestroyDequantizeGPTQDescriptor(infiniopDequantizeGPTQDescriptor_t desc);
 
 #endif

scripts/python_test.py

@@ -17,7 +17,7 @@ def run_tests(args):
         "causal_softmax.py",
         "clip.py",
         "conv.py",
-        "dequantize_awq.py",
+        # "dequantize_awq.py",
         "dequantize_gptq.py",
         "gelu.py",
         "gemm.py",

src/infiniop/ops/dequantize_awq/nvidia/dequantize_w42f16_kernel.cuh

@@ -122,4 +122,4 @@ __device__ uint4 dequantize_s4_to_fp16x2_awq(uint32_t const &source) {
     return result;
 #endif
     __builtin_unreachable(); // Suppress missing return statement warning
-}
+}

src/infiniop/ops/dequantize_gptq/moore/dequantize_w42f16_kernel.h

@@ -38,4 +38,4 @@ __device__ __forceinline__ uint4 dequantize_s4_to_fp16x2_gptq(uint32_t const &so
     result_ptr[2] = __halves2half2(hv2, hv6);
     result_ptr[3] = __halves2half2(hv3, hv7);
     return result;
-}
+}

src/infiniop/ops/dequantize_gptq/nvidia/dequantize_w42f16_kernel.cuh

@@ -17,7 +17,7 @@ __device__ uint4 dequantize_s4_to_fp16x2_gptq(uint32_t const &source) {
     // 步骤 2: GPTQ 是 (Q - Z) * S。
     // Q 和 Z 都是无符号数 [0, 15]。
     // 这里不需要 - offset
-    
+
     __half hv0 = __half(v0);
     __half hv1 = __half(v1);
     __half hv2 = __half(v2);
@@ -121,4 +121,4 @@ __device__ uint4 dequantize_s4_to_fp16x2_gptq(uint32_t const &source) {
     return result;
 #endif
     __builtin_unreachable(); // Suppress missing return statement warning
-}
+}

src/infiniop/ops/dequantize_gptq/nvidia/dequantize_w42f16_nvidia.cu

@@ -2,9 +2,9 @@
 
 #include "../../../devices/nvidia/nvidia_handle.cuh"
 #include "../../../devices/nvidia/nvidia_kernel_common.cuh"
+#include "../dequantize_gptq.h"
 #include "dequantize_w42f16_kernel.cuh"
 #include "dequantize_w42f16_nvidia.cuh"
-#include "../dequantize_gptq.h"
 #include <cuda_fp16.h>
 
 namespace op::dequantize_gptq::nvidia {
@@ -40,37 +40,41 @@ infiniStatus_t Descriptor::create(
 // zeros: [num_groups, out_packed] packing 8 output channels per word
 // scales: [num_groups, out_features], g_idx: [in_features]
 __global__ void __launch_bounds__(128)
-dequantize_weights_gptq(const uint32_t *__restrict__ qweight,
-                        const half     *__restrict__ scales,
-                        const uint32_t *__restrict__ zeros,
-                        const int      *__restrict__ g_idx,
-                        half           *__restrict__ out,
-                        int in_features,
-                        int out_features,
-                        int out_packed,   // ceil(out_features / 8)
-                        int num_groups) {
+    dequantize_weights_gptq(const uint32_t *__restrict__ qweight,
+                            const half *__restrict__ scales,
+                            const uint32_t *__restrict__ zeros,
+                            const int *__restrict__ g_idx,
+                            half *__restrict__ out,
+                            int in_features,
+                            int out_features,
+                            int out_packed, // ceil(out_features / 8)
+                            int num_groups) {
     // Each thread handles one packed output column (8 real output cols).
-    const int col_pack = blockIdx.x * blockDim.x + threadIdx.x;     // packed output column
-    const int row = blockIdx.y * blockDim.y + threadIdx.y;          // real input row
-    if (col_pack >= out_packed || row >= in_features) return;
+    const int col_pack = blockIdx.x * blockDim.x + threadIdx.x; // packed output column
+    const int row = blockIdx.y * blockDim.y + threadIdx.y;      // real input row
+    if (col_pack >= out_packed || row >= in_features) {
+        return;
+    }
 
     // Clamp gid to valid range
     const int gid_raw = g_idx ? g_idx[row] : 0;
     const int gid = ((gid_raw % num_groups) + num_groups) % num_groups;
 
-    const int pack_row = row >> 3;                                  // packed input row
+    const int pack_row = row >> 3; // packed input row
 
-    const int zero_idx = gid * out_packed + col_pack;               // zeros layout: [num_groups, out_packed]
+    const int zero_idx = gid * out_packed + col_pack; // zeros layout: [num_groups, out_packed]
     const uint32_t zeros_loaded = zeros[zero_idx];
 
-    const int q_shift = (row & 7) * 4;      // qweight packs 8 input rows
-    const int col_base = col_pack << 3;    // 8 real cols per pack
+    const int q_shift = (row & 7) * 4;  // qweight packs 8 input rows
+    const int col_base = col_pack << 3; // 8 real cols per pack
     const int scale_base = gid * out_features + col_base;
 
-    #pragma unroll
+#pragma unroll
     for (int j = 0; j < 8; ++j) {
         const int col = col_base + j;
-        if (col >= out_features) break;
+        if (col >= out_features) {
+            break;
+        }
 
         const uint32_t q_loaded = qweight[pack_row * out_features + col];
         const int q_nib = (q_loaded >> q_shift) & 0xF;
@@ -95,32 +99,33 @@ Descriptor::calculate(
     const void *g_idx,
     void *stream) const {
 
-    const int in_features  = _info.in_features();
+    const int in_features = _info.in_features();
     const int out_features = _info.out_features();
-    const int out_packed   = _info.out_packed();
-    const int in_packed    = _info.in_packed();
-    const int num_groups   = _info.num_groups();
+    const int out_packed = _info.out_packed();
+    const int in_packed = _info.in_packed();
+    const int num_groups = _info.num_groups();
 
-    if (num_groups <= 0 || in_features <= 0 || out_features <= 0 || out_packed <= 0 || in_packed <= 0)
+    if (num_groups <= 0 || in_features <= 0 || out_features <= 0 || out_packed <= 0 || in_packed <= 0) {
         return INFINI_STATUS_BAD_PARAM;
+    }
 
-    constexpr int BLOCK_X = 16;  // packed columns
-    constexpr int BLOCK_Y = 4;   // rows
+    constexpr int BLOCK_X = 16; // packed columns
+    constexpr int BLOCK_Y = 4;  // rows
     dim3 threads(BLOCK_X, BLOCK_Y);
     dim3 blocks((out_packed + BLOCK_X - 1) / BLOCK_X,
                 (in_features + BLOCK_Y - 1) / BLOCK_Y);
 
     dequantize_weights_gptq<<<blocks, threads, 0,
-        reinterpret_cast<cudaStream_t>(stream)>>>(
-            reinterpret_cast<const uint32_t*>(qweight),
-            reinterpret_cast<const half*>(scales),
-            reinterpret_cast<const uint32_t*>(zeros),
-            reinterpret_cast<const int*>(g_idx),
-            reinterpret_cast<half*>(out),
-            in_features, out_features, out_packed, num_groups);
+                              reinterpret_cast<cudaStream_t>(stream)>>>(
+        reinterpret_cast<const uint32_t *>(qweight),
+        reinterpret_cast<const half *>(scales),
+        reinterpret_cast<const uint32_t *>(zeros),
+        reinterpret_cast<const int *>(g_idx),
+        reinterpret_cast<half *>(out),
+        in_features, out_features, out_packed, num_groups);
     return INFINI_STATUS_SUCCESS;
 }
 
 } // namespace op::dequantize_gptq::nvidia
 
-#endif
+#endif

src/infiniop/ops/dequantize_gptq/operator.cc

@@ -12,22 +12,22 @@
 #include "iluvatar/dequantize_w42f16_iluvatar.cuh"
 #endif
 
-__C infiniStatus_t infiniopCreateDequantizeGPTQDescriptor(
+__INFINI_C infiniStatus_t infiniopCreateDequantizeGPTQDescriptor(
     infiniopHandle_t handle,
     infiniopDequantizeGPTQDescriptor_t *desc_ptr,
     infiniopTensorDescriptor_t out_desc,
     infiniopTensorDescriptor_t qweight_desc,
     infiniopTensorDescriptor_t scales_desc,
     infiniopTensorDescriptor_t zeros_desc,
-    infiniopTensorDescriptor_t g_idx_desc) {  // add g_idx
-#define CREATE(CASE, NAMESPACE)                                                       \
-    case CASE:                                                                        \
+    infiniopTensorDescriptor_t g_idx_desc) { // add g_idx
+#define CREATE(CASE, NAMESPACE)                                                        \
+    case CASE:                                                                         \
         return op::dequantize_gptq::NAMESPACE::Descriptor::create(                     \
-            handle,                                                                   \
+            handle,                                                                    \
             reinterpret_cast<op::dequantize_gptq::NAMESPACE::Descriptor **>(desc_ptr), \
-            out_desc,                                                                 \
-            qweight_desc,                                                             \
-            scales_desc,                                                              \
+            out_desc,                                                                  \
+            qweight_desc,                                                              \
+            scales_desc,                                                               \
             zeros_desc, g_idx_desc)
 
     switch (handle->device) {
@@ -50,10 +50,10 @@ __C infiniStatus_t infiniopCreateDequantizeGPTQDescriptor(
 #undef CREATE
 }
 
-__C infiniStatus_t infiniopGetDequantizeGPTQWorkspaceSize(infiniopDequantizeGPTQDescriptor_t desc,
-                                                         size_t *size) {
-#define GET(CASE, NAMESPACE)                                                                                \
-    case CASE:                                                                                              \
+__INFINI_C infiniStatus_t infiniopGetDequantizeGPTQWorkspaceSize(infiniopDequantizeGPTQDescriptor_t desc,
+                                                                 size_t *size) {
+#define GET(CASE, NAMESPACE)                                                                                 \
+    case CASE:                                                                                               \
         *size = reinterpret_cast<const op::dequantize_gptq::NAMESPACE::Descriptor *>(desc)->workspaceSize(); \
         return INFINI_STATUS_SUCCESS
 
@@ -76,19 +76,19 @@ __C infiniStatus_t infiniopGetDequantizeGPTQWorkspaceSize(infiniopDequantizeGPTQ
 #undef GET
 }
 
-__C infiniStatus_t infiniopDequantizeGPTQ(
+__INFINI_C infiniStatus_t infiniopDequantizeGPTQ(
     infiniopDequantizeGPTQDescriptor_t desc,
     void *workspace,
     size_t workspace_size,
     void *out,
     const void *qweight,
     const void *scales,
     const void *zeros,
-    const void *g_idx,    // add g_idx
+    const void *g_idx, // add g_idx
     void *stream) {
 
-#define CALCULATE(CASE, NAMESPACE)                                                       \
-    case CASE:                                                                           \
+#define CALCULATE(CASE, NAMESPACE)                                                        \
+    case CASE:                                                                            \
         return reinterpret_cast<const op::dequantize_gptq::NAMESPACE::Descriptor *>(desc) \
             ->calculate(workspace, workspace_size, out, qweight, scales, zeros, g_idx, stream)
 
@@ -112,11 +112,11 @@ __C infiniStatus_t infiniopDequantizeGPTQ(
 #undef CALCULATE
 }
 
-__C infiniStatus_t
+__INFINI_C infiniStatus_t
 infiniopDestroyDequantizeGPTQDescriptor(infiniopDequantizeGPTQDescriptor_t desc) {
 
-#define DELETE(CASE, NAMESPACE)                                                           \
-    case CASE:                                                                            \
+#define DELETE(CASE, NAMESPACE)                                                            \
+    case CASE:                                                                             \
         delete reinterpret_cast<const op::dequantize_gptq::NAMESPACE::Descriptor *>(desc); \
         return INFINI_STATUS_SUCCESS;
 
@@ -140,4 +140,4 @@ infiniopDestroyDequantizeGPTQDescriptor(infiniopDequantizeGPTQDescriptor_t desc)
 #undef DELETE
 }
 
-// #endif
+// #endif