diff --git a/src/blas/backends/cublas/cublas_batch.cpp b/src/blas/backends/cublas/cublas_batch.cpp
index 6fd32a62e..911f8b5e8 100644
--- a/src/blas/backends/cublas/cublas_batch.cpp
+++ b/src/blas/backends/cublas/cublas_batch.cpp
@@ -16,7 +16,9 @@
 *  limitations under the License.
 *
 **************************************************************************/
+# ifndef __HIPSYCL__
 #include <CL/sycl/detail/pi.hpp>
+# endif
 #include "cublas_helper.hpp"
 #include "cublas_scope_handle.hpp"
 #include "oneapi/mkl/exceptions.hpp"
diff --git a/src/blas/backends/cublas/cublas_extensions.cpp b/src/blas/backends/cublas/cublas_extensions.cpp
index 3df6c8360..d44115abf 100644
--- a/src/blas/backends/cublas/cublas_extensions.cpp
+++ b/src/blas/backends/cublas/cublas_extensions.cpp
@@ -16,7 +16,9 @@
 *  limitations under the License.
 *
 **************************************************************************/
+#ifndef __HIPSYCL__
 #include <CL/sycl/detail/pi.hpp>
+#endif
 #include "cublas_helper.hpp"
 #include "cublas_scope_handle.hpp"
 #include "oneapi/mkl/exceptions.hpp"
diff --git a/src/blas/backends/cublas/cublas_level1.cpp b/src/blas/backends/cublas/cublas_level1.cpp
index 0ff69e5d1..2e259852e 100644
--- a/src/blas/backends/cublas/cublas_level1.cpp
+++ b/src/blas/backends/cublas/cublas_level1.cpp
@@ -20,9 +20,9 @@
 #include "cublas_scope_handle.hpp"
 #include "oneapi/mkl/exceptions.hpp"
 #include "oneapi/mkl/blas/detail/cublas/onemkl_blas_cublas.hpp"
-
+#ifndef __HIPSYCL__
 #include <CL/sycl/detail/pi.hpp>
-
+#endif
 namespace oneapi {
 namespace mkl {
 namespace blas {
diff --git a/src/blas/backends/cublas/cublas_level2.cpp b/src/blas/backends/cublas/cublas_level2.cpp
index db23a3eb6..e3c98b749 100644
--- a/src/blas/backends/cublas/cublas_level2.cpp
+++ b/src/blas/backends/cublas/cublas_level2.cpp
@@ -16,7 +16,10 @@
 *  limitations under the License.
 *
 **************************************************************************/
+#ifndef __HIPSYCL__
 #include <CL/sycl/detail/pi.hpp>
+#endif
+
 #include "cublas_helper.hpp"
 #include "cublas_scope_handle.hpp"
 #include "oneapi/mkl/exceptions.hpp"
diff --git a/src/blas/backends/cublas/cublas_level3.cpp b/src/blas/backends/cublas/cublas_level3.cpp
index 671a15ea7..854a26929 100644
--- a/src/blas/backends/cublas/cublas_level3.cpp
+++ b/src/blas/backends/cublas/cublas_level3.cpp
@@ -16,7 +16,9 @@
 *  limitations under the License.
 *
 **************************************************************************/
+#ifndef __HIPSYCL__
 #include <CL/sycl/detail/pi.hpp>
+#endif
 #include "cublas_helper.hpp"
 #include "cublas_scope_handle.hpp"
 #include "oneapi/mkl/exceptions.hpp"
diff --git a/src/blas/backends/cublas/cublas_scope_handle.cpp b/src/blas/backends/cublas/cublas_scope_handle.cpp
index b0ef3b378..0d28ddc51 100644
--- a/src/blas/backends/cublas/cublas_scope_handle.cpp
+++ b/src/blas/backends/cublas/cublas_scope_handle.cpp
@@ -17,13 +17,17 @@
 *
 **************************************************************************/
 #include "cublas_scope_handle.hpp"
+#ifndef __HIPSYCL__
 #include <CL/sycl/detail/common.hpp>
+#include <cuda.h>
+#endif
 
 namespace oneapi {
 namespace mkl {
 namespace blas {
 namespace cublas {
 
+
 cublas_handle::~cublas_handle() noexcept(false) {
     for (auto &handle_pair : cublas_handle_mapper_) {
         cublasStatus_t err;
@@ -39,6 +43,7 @@ cublas_handle::~cublas_handle() noexcept(false) {
     }
     cublas_handle_mapper_.clear();
 }
+
 /**
  * Inserts a new element in the map if its key is unique. This new element
  * is constructed in place using args as the arguments for the construction
@@ -47,11 +52,14 @@ cublas_handle::~cublas_handle() noexcept(false) {
  * the one being emplaced (keys in a map container are unique).
  */
 thread_local cublas_handle CublasScopedContextHandler::handle_helper = cublas_handle{};
-
+#ifdef __HIPSYCL__
+CublasScopedContextHandler::CublasScopedContextHandler(cl::sycl::queue queue, cl::sycl::interop_handle ih) : interop_h(ih){}
+#else
 CublasScopedContextHandler::CublasScopedContextHandler(cl::sycl::queue queue) {
     placedContext_ = queue.get_context();
     auto device = queue.get_device();
     auto desired = cl::sycl::get_native<cl::sycl::backend::cuda>(placedContext_);
+
     auto cudaDevice = cl::sycl::get_native<cl::sycl::backend::cuda>(device);
     CUresult err;
     CUDA_ERROR_FUNC(cuCtxGetCurrent, err, &original_);
@@ -71,12 +79,15 @@ CublasScopedContextHandler::CublasScopedContextHandler(cl::sycl::queue queue) {
         needToRecover_ = !(original_ == nullptr && isPrimary);
     }
 }
+#endif
 
 CublasScopedContextHandler::~CublasScopedContextHandler() noexcept(false) {
+    #ifndef __HIPSYCL__  
     if (needToRecover_) {
         CUresult err;
         CUDA_ERROR_FUNC(cuCtxSetCurrent, err, original_);
     }
+    #endif
 }
 
 void ContextCallback(void *userData) {
@@ -97,10 +108,19 @@ void ContextCallback(void *userData) {
 }
 
 cublasHandle_t CublasScopedContextHandler::get_handle(const cl::sycl::queue &queue) {
+    #ifndef __HIPSYCL__
     auto piPlacedContext_ = reinterpret_cast<pi_context>(placedContext_.get());
+    #else
+    cl::sycl::device device = queue.get_device();
+    int current_device = interop_h.get_native_device<cl::sycl::backend::cuda>();
+    #endif
     CUstream streamId = get_stream(queue);
     cublasStatus_t err;
+    #ifndef __HIPSYCL__
     auto it = handle_helper.cublas_handle_mapper_.find(piPlacedContext_);
+    #else
+    auto it = handle_helper.cublas_handle_mapper_.find(current_device);
+    #endif
     if (it != handle_helper.cublas_handle_mapper_.end()) {
         if (it->second == nullptr) {
             handle_helper.cublas_handle_mapper_.erase(it);
@@ -120,29 +140,36 @@ cublasHandle_t CublasScopedContextHandler::get_handle(const cl::sycl::queue &que
             }
         }
     }
-
     cublasHandle_t handle;
 
     CUBLAS_ERROR_FUNC(cublasCreate, err, &handle);
     CUBLAS_ERROR_FUNC(cublasSetStream, err, handle, streamId);
 
     auto insert_iter = handle_helper.cublas_handle_mapper_.insert(
+        #ifdef __HIPSYCL__
+        std::make_pair(current_device, new std::atomic<cublasHandle_t>(handle)));
+        #else
         std::make_pair(piPlacedContext_, new std::atomic<cublasHandle_t>(handle)));
-
+        #endif
+    #ifndef __HIPSYCL__
     auto ptr = &(insert_iter.first->second);
 
     sycl::detail::pi::contextSetExtendedDeleter(placedContext_, ContextCallback, ptr);
-
+    #endif
     return handle;
 }
-
+#ifndef __HIPSYCL__
 CUstream CublasScopedContextHandler::get_stream(const cl::sycl::queue &queue) {
     return cl::sycl::get_native<cl::sycl::backend::cuda>(queue);
 }
 cl::sycl::context CublasScopedContextHandler::get_context(const cl::sycl::queue &queue) {
     return queue.get_context();
 }
-
+#else
+CUstream CublasScopedContextHandler::get_stream(const cl::sycl::queue &queue) {
+    return interop_h.get_native_queue<cl::sycl::backend::cuda>();
+}
+#endif
 } // namespace cublas
 } // namespace blas
 } // namespace mkl
diff --git a/src/blas/backends/cublas/cublas_scope_handle.hpp b/src/blas/backends/cublas/cublas_scope_handle.hpp
index e73c63225..3b7bf105b 100644
--- a/src/blas/backends/cublas/cublas_scope_handle.hpp
+++ b/src/blas/backends/cublas/cublas_scope_handle.hpp
@@ -19,9 +19,11 @@
 #ifndef _MKL_BLAS_CUBLAS_SCOPED_HANDLE_HPP_
 #define _MKL_BLAS_CUBLAS_SCOPED_HANDLE_HPP_
 #include <CL/sycl.hpp>
+#ifndef __HIPSYCL__
 #include <CL/sycl/backend/cuda.hpp>
 #include <CL/sycl/context.hpp>
 #include <CL/sycl/detail/pi.hpp>
+#endif
 #include <atomic>
 #include <memory>
 #include <thread>
@@ -32,11 +34,19 @@ namespace mkl {
 namespace blas {
 namespace cublas {
 
+#ifndef __HIPSYCL__
 struct cublas_handle {
     using handle_container_t = std::unordered_map<pi_context, std::atomic<cublasHandle_t> *>;
     handle_container_t cublas_handle_mapper_{};
     ~cublas_handle() noexcept(false);
 };
+#else
+struct cublas_handle {
+    using handle_container_t = std::unordered_map<int, std::atomic<cublasHandle_t>* >;
+    handle_container_t cublas_handle_mapper_{};
+    ~cublas_handle() noexcept(false);
+};
+#endif
 
 /**
 * @brief NVIDIA advise for handle creation:
@@ -48,7 +58,6 @@ According to NVIDIA:
 http://docs.nvidia.com/cuda/cublas/index.html#thread-safety2changeme
 3)	It is neither required nor recommended that different handles be used for different streams on the same device,
  using the same host thread.
-
 However, the 3 above advises are for using cuda runtime API. The NVIDIA runtime API creates a default context for users. 
 The createHandle function in cuBLAS uses the context located on top of the stack for each thread. Then, the cuBLAS routine 
 uses this context for resource allocation/access. Calling a cuBLAS function with a handle created for context A and 
@@ -65,16 +74,23 @@ the handle must be destroyed when the context goes out of scope. This will bind
 **/
 
 class CublasScopedContextHandler {
+    #ifdef __HIPSYCL__
+    cl::sycl::interop_handle interop_h;
+    #else
     CUcontext original_;
     cl::sycl::context placedContext_;
     bool needToRecover_;
+    #endif
     static thread_local cublas_handle handle_helper;
+
     CUstream get_stream(const cl::sycl::queue &queue);
     cl::sycl::context get_context(const cl::sycl::queue &queue);
 
+
 public:
     CublasScopedContextHandler(cl::sycl::queue queue);
-
+    explicit CublasScopedContextHandler(cl::sycl::queue queue, cl::sycl::interop_handle ih);
+    
     ~CublasScopedContextHandler() noexcept(false);
     /**
    * @brief get_handle: creates the handle by implicitly impose the advice
@@ -86,11 +102,18 @@ class CublasScopedContextHandler {
     cublasHandle_t get_handle(const cl::sycl::queue &queue);
     // This is a work-around function for reinterpret_casting the memory. This
     // will be fixed when SYCL-2020 has been implemented for Pi backend.
+    #ifdef __HIPSYCL__
     template <typename T, typename U>
+    inline T get_mem(cl::sycl::interop_handle ih, U acc) {
+        return reinterpret_cast<T>(ih.get_native_mem<cl::sycl::backend::cuda>(acc));
+    }
+    #else
+    template <typename T, typename U >
     inline T get_mem(cl::sycl::interop_handler ih, U acc) {
         CUdeviceptr cudaPtr = ih.get_mem<cl::sycl::backend::cuda>(acc);
         return reinterpret_cast<T>(cudaPtr);
     }
+    #endif
 };
 
 } // namespace cublas