[GCU] Remove unused checkGcuErrors and restore axpy operator logic

- Remove unused checkGcuErrors macro and __checkGcuErrors function
  from jit_utils.h; GCU error handling is centralized in gcu_backend.h
- Remove unnecessary #include "tops_runtime_api.h" from jit_utils.h
- Restore axpy3_kernel to original `o *= a` form, reverting the
  workaround that is no longer needed after fp64-to-fp32 signature
  conversion was implemented in standalone_compile.py

Co-authored-by: Cursor <cursoragent@cursor.com>

Author: baoqiliu

examples/arg_handle/axpy.py

@@ -41,7 +41,7 @@ def axpy3_kernel(X, Y, Out, a, n, BLOCK_N: tl.constexpr):
     x = tl.load(X + offsets, mask=mask)
     o = x
     if a is not None:
-        o = o * a
+        o *= a
 
     if Y is not None:
         y = tl.load(Y + offsets, mask=mask)

examples/arg_handle/test_axpy.cpp

@@ -2,6 +2,7 @@
 #include "axpy_op.h"
 #include "torch/torch.h"
 #include "triton_jit/backend_config.h"
+
 static at::Device test_device() {
 #if defined(BACKEND_NPU)
   return at::Device("npu:0");

include/triton_jit/jit_utils.h

@@ -17,7 +17,6 @@
 #elif defined(BACKEND_MACA)
 #include <mcr/mc_runtime.h>
 #elif defined(BACKEND_GCU)
-#include "tops_runtime_api.h"
 #else
 #include "cuda.h"
 #endif
@@ -169,20 +168,7 @@ inline void __checkMacaErrors(mcError_t code, const char* file, const int line)
   }
 }
 #elif defined(BACKEND_GCU)
-#define checkGcuErrors(err) __checkGcuErrors(err, __FILE__, __LINE__)
-
-inline void __checkGcuErrors(topsError_t code, const char* file, const int line) {
-  if (code != topsSuccess) {
-    const char* error_string = topsGetErrorString(code);
-    fprintf(stderr,
-            "GCU TOPS API error = %04d from file <%s>, line %i. Detail: <%s>\n",
-            static_cast<int>(code),
-            file,
-            line,
-            error_string);
-    throw std::runtime_error(error_string ? error_string : "Unknown GCU error");
-  }
-}
+// GCU error handling is done inline in gcu_backend.h
 #else
 void ensure_cuda_context();