[Bugfix] Make quickstart work properly on cu118 (#193)

src/runtime/runtime.cc

@@ -17,6 +17,7 @@ namespace tl {
 
 using namespace runtime;
 
+#if (__CUDACC_VER_MAJOR__ >= 12)
 template <typename T> static std::string ArrayToStr(const T *ptr, size_t n) {
   std::stringstream ss;
   ss << "[";
@@ -202,6 +203,7 @@ TVM_REGISTER_GLOBAL(tvm_tensormap_create_im2col)
       }
       *ret = static_cast<int>(result);
     });
+#endif // (__CUDACC_VER_MAJOR__ >= 12)
 
 } // namespace tl
 } // namespace tvm

src/runtime/runtime.h

@@ -13,10 +13,12 @@
 namespace tvm {
 namespace tl {
 
+#if (__CUDACC_VER_MAJOR__ >= 12)
 constexpr const char *tvm_tensormap_create_tiled =
     "__tvm_tensormap_create_tiled";
 constexpr const char *tvm_tensormap_create_im2col =
     "__tvm_tensormap_create_im2col";
+#endif // (__CUDACC_VER_MAJOR__ >= 12)
 } // namespace tl
 } // namespace tvm
 

src/transform/lower_hopper_intrin.cc

@@ -36,6 +36,7 @@ namespace tl {
 
 using namespace tir;
 
+#if (__CUDACC_VER_MAJOR__ >= 12)
 class LowerHopperIntrin : public StmtExprMutator {
 public:
   static PrimFunc Substitute(PrimFunc &f) {
@@ -168,6 +169,7 @@ tvm::transform::Pass LowerHopperIntrin() {
 
 TVM_REGISTER_GLOBAL("tl.transform.LowerHopperIntrin")
     .set_body_typed(LowerHopperIntrin);
+#endif // (__CUDACC_VER_MAJOR__ >= 12)
 
 } // namespace tl
 } // namespace tvm

tilelang/transform/__init__.py

@@ -81,7 +81,8 @@ def LowerHopperIntrin():
     fpass : tvm.transform.Pass
         The result pass
     """
-    return _ffi_api.LowerHopperIntrin()  # type: ignore
+    return _ffi_api.LowerHopperIntrin() \
+        if hasattr(_ffi_api, "LowerHopperIntrin") else lambda f: f # type: ignore
 
 
 def WarpSpecializedPipeline():