NVIDIA · zcsizmadia · May 8, 2026 · May 8, 2026 · May 8, 2026 · May 8, 2026
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -152,6 +152,68 @@ rapids_cpm_cccl(
 
 target_link_libraries(matx INTERFACE CCCL::CCCL)
 
+if(WIN32)
+    # Prevent wingdi.h from defining ERROR=0, which breaks MATX_LOG_ERROR
+    target_compile_definitions(matx INTERFACE NOMINMAX WIN32_LEAN_AND_MEAN NOGDI _USE_MATH_DEFINES)
+
+    if(MSVC)
+        # /Zc:preprocessor: MSVC conformance preprocessor required by libcudacxx / MatX.
+        # Passed as -Xcompiler so it reaches cl.exe through nvcc.
+        add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/Zc:preprocessor>")
+        add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:/Zc:preprocessor>")
+
+        # /Zc:__cplusplus: Without this MSVC always reports __cplusplus=199711L regardless
+        # of /std:c++NN, which breaks any library (e.g. nvbench) that checks __cplusplus
+        # directly in .cxx files compiled by cl.exe (not through nvcc). The equivalent
+        # -Xcompiler=/Zc:__cplusplus in CMAKE_CUDA_FLAGS only reaches cl.exe for .cu files.
+        add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/Zc:__cplusplus>")
+        add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:/Zc:__cplusplus>")
+
+        # /bigobj raises the COFF section limit from 2^16 to 2^32,
+        # required for heavily-templated translation units.
+        add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/bigobj>")
+        add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:/bigobj>")
+
+        # /wd4996: suppress "deprecated" warnings from NVIDIA's own headers (cusparse.h,
+        # cusolverDn.h, etc.) that annotate types scheduled for removal. These warnings
+        # originate in external SDK headers and cannot be fixed in MatX; without /wd4996
+        # they become errors under /WX.
+        add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/wd4996>")
+
+        # /wd4100: suppress unreferenced formal parameter warnings from third-party headers
+        # (Python 3.14 cpython/unicodeobject.h has intentional _unused_op parameters).
+        add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/wd4100>")
+        add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:/wd4100>")
+
+        # /wd4864: VS18 new warning for dependent-name template calls on free function
+        # templates. Correct code, not a real ambiguity; suppress to unblock /WX builds.
+        add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/wd4864>")
+        add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:/wd4864>")
+
+        # /wd4702: unreachable code in constexpr-if chains — MSVC false positive when every
+        # if-constexpr branch returns but a fallback return is still needed for non-MSVC.
+        add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/wd4702>")
+        add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:/wd4702>")
+
+        # /wd4324: structure padded due to alignas — intentional in CUDA operator types.
+        add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/wd4324>")
+        add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:/wd4324>")
+
+        # /wd4714: __forceinline function not inlined — expected for large templated
+        # functions such as get_type_str() instantiated over complex operator types.
+        add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/wd4714>")
+        add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:/wd4714>")
+
+        # /wd4127: conditional expression is constant — common in template/constexpr code.
+        # VS18 2025+ (MSVC 1950+) resolves most of these automatically; suppress only on
+        # VS17 2022 and earlier.
+        if(MSVC_VERSION LESS 1950)
+            add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/wd4127>")
+            add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:/wd4127>")
+        endif()
+    endif()
+endif()
+
 # Set flags for compiling tests faster (only for nvcc)
 if (NOT CMAKE_CUDA_COMPILER_ID STREQUAL "Clang")
     set(MATX_CUDA_FLAGS ${CMAKE_CUDA_FLAGS} --threads 0 -ftemplate-backtrace-limit=0 --extended-lambda)

diff --git a/bench/00_operators/operators.cu b/bench/00_operators/operators.cu
@@ -132,7 +132,9 @@ void sphericalharmonics(nvbench::state &state, nvbench::type_list<ValueType>)
 
   state.exec( 
     [&](nvbench::launch &launch) {
+MATX_IGNORE_WARNING_PUSH_MSVC(4834) // expression result unused
       (X=XXm, Y=YYm, Z=ZZm).run((cudaStream_t)launch.get_stream());
+MATX_IGNORE_WARNING_POP_MSVC
     });
 }
 

diff --git a/bench/00_transform/sarbp.cu b/bench/00_transform/sarbp.cu
@@ -37,7 +37,7 @@ using namespace matx;
 
 // SAR backprojection benchmarks for different precision modes
 
-const std::vector<ssize_t> PROBLEM_SIZES = {2000};
+const std::vector<int64_t> PROBLEM_SIZES = {2000};
 
 // Gigabackprojections per second derived metric. Operations per launch =
 // num_pulses * image_width * image_height = problem_size^3. Reading this

diff --git a/bench/CMakeLists.txt b/bench/CMakeLists.txt
@@ -34,7 +34,20 @@ function(create_bench_executable bench_file)
     set_property(TARGET ${target_name} PROPERTY ENABLE_EXPORTS 1)
 
     if (MSVC)
-        target_compile_options(${target_name} PRIVATE /W4 /WX)
+        # /W4 and /WX must be wrapped in -Xcompiler for CUDA source files.
+        # nvcc does not recognize MSVC-style /flags as its own options; passing them
+        # directly without -Xcompiler causes nvcc to treat them as extra input files,
+        # triggering "A single input file is required for a non-link phase".
+        target_compile_options(${target_name} PRIVATE
+            $<$<COMPILE_LANGUAGE:CXX>:/W4>
+            $<$<COMPILE_LANGUAGE:CXX>:/WX>
+            $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/W4>
+            $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/WX>
+        )
+
+        # cudart_static.lib pulls in LIBCMT (static CRT); suppress the LNK4098 warning
+        # that results from mixing it with the /MD (dynamic CRT) runtime used elsewhere.
+        target_link_options(${target_name} PRIVATE /NODEFAULTLIB:LIBCMT)
     else()
         target_compile_options(${target_name} PRIVATE ${WARN_FLAGS})
         target_compile_options(${target_name} PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:${MATX_CUDA_FLAGS}>)

diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
@@ -39,7 +39,20 @@ if(eigen_DIR)
 endif()
 
 if (MSVC)
-    target_compile_options(example_lib INTERFACE /W4 /WX)
+    # /W4 and /WX must be wrapped in -Xcompiler for CUDA source files.
+    # nvcc does not recognize MSVC-style /flags as its own options; passing them
+    # directly without -Xcompiler causes nvcc to treat them as extra input files,
+    # triggering "A single input file is required for a non-link phase".
+    target_compile_options(example_lib INTERFACE
+        $<$<COMPILE_LANGUAGE:CXX>:/W4>
+        $<$<COMPILE_LANGUAGE:CXX>:/WX>
+        $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/W4>
+        $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/WX>
+    )
+
+    # cudart_static.lib pulls in LIBCMT (static CRT); suppress the LNK4098 warning
+    # that results from mixing it with the /MD (dynamic CRT) runtime used elsewhere.
+    target_link_options(example_lib INTERFACE /NODEFAULTLIB:LIBCMT)
 else()
     target_compile_options(example_lib INTERFACE ${WARN_FLAGS})
     target_compile_options(example_lib INTERFACE ${MATX_CUDA_FLAGS})
@@ -82,7 +95,15 @@ endforeach()
 # Add host-compiler only example program to catch missing ifdef __CUDACC__ guards
 add_executable(test_host test_host.cpp)
 target_link_libraries(test_host matx::matx)
-target_compile_options(test_host PRIVATE -Wall -Werror)
+if (MSVC)
+    # -Wall/-Werror are GCC/Clang flags; cl.exe uses /WX for warnings-as-errors.
+    # /Werror is not a valid MSVC option and causes D8021.
+    target_compile_options(test_host PRIVATE /WX)
+    # matx::matx transitively links cudart_static which pulls in LIBCMT; suppress LNK4098.
+    target_link_options(test_host PRIVATE /NODEFAULTLIB:LIBCMT)
+else()
+    target_compile_options(test_host PRIVATE -Wall -Werror)
+endif()
 
 if(MATX_EN_PYBIND11)
     add_subdirectory(python_integration_sample)

diff --git a/examples/python_integration_sample/CMakeLists.txt b/examples/python_integration_sample/CMakeLists.txt
@@ -53,6 +53,13 @@ endif()
 
 add_library(matxutil MODULE matxutil.cu)
 target_link_libraries(matxutil PRIVATE matx::matx CUDA::cudart)
+
+if (MSVC)
+    # cudart_static.lib pulls in LIBCMT (static CRT); suppress the LNK4098 warning
+    # that results from mixing it with the /MD (dynamic CRT) runtime used elsewhere.
+    target_link_options(matxutil PRIVATE /NODEFAULTLIB:LIBCMT)
+endif()
+
 set_target_properties(matxutil PROPERTIES SUFFIX ".so" PREFIX "")
 
 configure_file(

diff --git a/examples/spectrogram_graph.cu b/examples/spectrogram_graph.cu
@@ -57,7 +57,7 @@ int main([[maybe_unused]] int argc, [[maybe_unused]] char **argv)
 
   using complex = cuda::std::complex<float>;
   cudaGraph_t graph;
-  cudaGraphExec_t instance;
+  cudaGraphExec_t instance = nullptr;
 
   cudaStream_t stream;
   cudaStreamCreate(&stream);

diff --git a/examples/spherical_harmonics.cu b/examples/spherical_harmonics.cu
@@ -78,7 +78,9 @@ int main([[maybe_unused]] int argc, [[maybe_unused]] char **argv)
   auto Y = make_tensor<ValueType>(Ym.Shape());
   auto Z = make_tensor<ValueType>(Zm.Shape());
 
+MATX_IGNORE_WARNING_PUSH_MSVC(4834)
   (X = Xm, Y = Ym, Z=Zm).run(exec);
+MATX_IGNORE_WARNING_POP_MSVC
 
   exec.sync();
 

diff --git a/include/matx/core/allocator.h b/include/matx/core/allocator.h
@@ -290,7 +290,7 @@ struct MemTracker {
 
 
 
-__attribute__ ((visibility ("default")))
+MATX_VISIBILITY_DEFAULT
 __MATX_INLINE__ MemTracker &GetAllocMap() {
   static MemTracker tracker;
   return tracker;
@@ -304,7 +304,7 @@ __MATX_INLINE__ MemTracker &GetAllocMap() {
 // call matxFree() on user-managed pointers after calling this function. This function should
 // be called after the user application has called matxFree() on any pointers for which it
 // will call matxFree().
-__attribute__ ((visibility ("default")))
+MATX_VISIBILITY_DEFAULT
 __MATX_INLINE__ void FreeAllocations() {
   GetAllocMap().free_all();
 }
@@ -389,7 +389,7 @@ __MATX_INLINE__ void matxPrintMemoryStatistics()
   matxGetMemoryStats(&current, &total, &max);
 
   printf("Memory Statistics(GB):  current: %.2f, total: %.2f, max: %.2f. Total "
-         "allocations: %lu\n",
+         "allocations: %zu\n",
          static_cast<double>(current) / 1e9, static_cast<double>(total) / 1e9,
          static_cast<double>(max) / 1e9, GetAllocMap().size());
 }

diff --git a/include/matx/core/cache.h b/include/matx/core/cache.h
@@ -116,7 +116,7 @@ struct CacheCommonParamsKeyHash {
 };
 
 #ifndef DOXYGEN_ONLY
-__attribute__ ((visibility ("default")))
+MATX_VISIBILITY_DEFAULT
 #endif
 inline cuda::std::atomic<CacheId> CacheIdCounter{0};
 inline std::recursive_mutex cache_mtx; ///< Mutex protecting updates from map
@@ -130,7 +130,7 @@ inline auto& CacheRegistry() {
 }
 
 template<typename CacheType>
-__attribute__ ((visibility ("default")))
+MATX_VISIBILITY_DEFAULT
 CacheId GetCacheIdFromType()
 {
   static CacheId id = CacheIdCounter.fetch_add(1);
@@ -732,13 +732,13 @@ inline size_t PodArrayToHash(cuda::std::array<T, len> c)
   return hash;
 }
 
-__attribute__ ((visibility ("default")))
+MATX_VISIBILITY_DEFAULT
 __MATX_INLINE__ matxCache_t &InitCache() {
   static matxCache_t cache;
   return cache;
 }
 
-__attribute__ ((visibility ("default")))
+MATX_VISIBILITY_DEFAULT
 __MATX_INLINE__ matxCache_t &GetCache() {
   [[maybe_unused]] const auto &tracker = GetAllocMap();
   return InitCache();
@@ -751,15 +751,15 @@ __MATX_INLINE__ matxCache_t &GetCache() {
 // function does not clear the allocator cache (i.e., allocations made with matxAlloc
 // other than those created to support transforms).
 // To free the allocator cache, use matx::FreeAllocations().
-__attribute__ ((visibility ("default")))
+MATX_VISIBILITY_DEFAULT
 __MATX_INLINE__ void ClearCaches() {
   detail::GetCache().ClearAll();
 }
 
 // Helper function to clear both MatX caches and allocations. This provides a single
 // function that can be called prior to program exit to support clean shutdown
 // (i.e., to avoid issues with the order of destruction of static objects and CUDA contexts).
-__attribute__ ((visibility ("default")))
+MATX_VISIBILITY_DEFAULT
 __MATX_INLINE__ void ClearCachesAndAllocations() {
   ClearCaches();
   FreeAllocations();

diff --git a/include/matx/core/capabilities.h b/include/matx/core/capabilities.h
@@ -360,7 +360,7 @@ namespace detail {
     // ( (static_assert(std::is_same_v<ChildCapTypes, CapType>,"Child capability type mismatch.")), ...); // C++17 fold for static_assert
 
     CapabilityQueryType query_type = get_query_type(Cap);
-    CapType children_aggregated_val;
+    CapType children_aggregated_val{};
 
     // Step 1: Aggregate children capabilities
     if constexpr (sizeof...(ChildCapTypes) == 0) {

diff --git a/include/matx/core/defines.h b/include/matx/core/defines.h
@@ -80,6 +80,9 @@ namespace matx {
 #if defined(__clang__ )
     #define MATX_IGNORE_WARNING_PUSH_GCC(WARN_MSG)
     #define MATX_IGNORE_WARNING_POP_GCC
+    #define MATX_IGNORE_WARNING_PUSH_MSVC(WARN_ID)
+    #define MATX_IGNORE_WARNING_POP_MSVC
+    #define MATX_IGNORE_WARNING_NEXT_LINE_MSVC(WARN_ID)
 
     #define MATX_IGNORE_WARNING_PUSH_CLANG(WARN_MSG) \
         _Pragma("clang diagnostic push") \
@@ -90,24 +93,58 @@ namespace matx {
 #elif defined(__GNUC__)
     #define MATX_IGNORE_WARNING_PUSH_CLANG(WARN_MSG)
     #define MATX_IGNORE_WARNING_POP_CLANG
+    #define MATX_IGNORE_WARNING_PUSH_MSVC(WARN_ID)
+    #define MATX_IGNORE_WARNING_POP_MSVC
+    #define MATX_IGNORE_WARNING_NEXT_LINE_MSVC(WARN_ID)
 
     #define MATX_IGNORE_WARNING_PUSH_GCC(WARN_MSG) \
         _Pragma("GCC diagnostic push") \
         _Pragma(MATX_TOSTRING(GCC diagnostic ignored WARN_MSG))
 
     #define MATX_IGNORE_WARNING_POP_GCC \
         _Pragma("GCC diagnostic pop")
+#elif defined(_MSC_VER)
+    #define MATX_IGNORE_WARNING_PUSH_GCC(WARN_MSG)
+    #define MATX_IGNORE_WARNING_POP_GCC
+    #define MATX_IGNORE_WARNING_PUSH_CLANG(WARN_MSG)
+    #define MATX_IGNORE_WARNING_POP_CLANG
+
+    #define MATX_IGNORE_WARNING_PUSH_MSVC(WARN_ID) \
+        __pragma(warning(push)) \
+        __pragma(warning(disable : WARN_ID))
+
+    #define MATX_IGNORE_WARNING_POP_MSVC \
+        __pragma(warning(pop))
+
+    #define MATX_IGNORE_WARNING_NEXT_LINE_MSVC(WARN_ID) \
+        __pragma(warning(suppress : WARN_ID))
 #else
     #define MATX_IGNORE_WARNING_PUSH_GCC(WARN_MSG)
     #define MATX_IGNORE_WARNING_POP_GCC
     #define MATX_IGNORE_WARNING_PUSH_CLANG(WARN_MSG)
     #define MATX_IGNORE_WARNING_POP_CLANG
+    #define MATX_IGNORE_WARNING_PUSH_MSVC(WARN_ID)
+    #define MATX_IGNORE_WARNING_POP_MSVC
+    #define MATX_IGNORE_WARNING_NEXT_LINE_MSVC(WARN_ID)
 #endif
 
 // std::ceil is not constexpr until C++23
 #define MATX_ROUND_UP(N, S) ((((N) + (S) - 1) / (S)) * (S))
 
-enum {
+// MATX_VISIBILITY_DEFAULT: marks a symbol for export in shared library builds.
+// On GCC/Clang this expands to __attribute__((visibility("default"))).
+// On MSVC, __attribute__ is not supported and visibility is controlled by
+// __declspec(dllexport/dllimport), so the attribute is silenced here.
+#if defined(_MSC_VER)
+  #define MATX_VISIBILITY_DEFAULT
+#else
+  #define MATX_VISIBILITY_DEFAULT __attribute__ ((visibility ("default")))
+#endif
+
+// Explicit underlying type prevents MSVC C4309 ("truncation of constant
+// value") when index_t is 64-bit: a plain enum defaults to int (32-bit),
+// causing LLONG_MAX to truncate and /WX to turn that into an error.
+enum : index_t {
   matxKeepDim     = cuda::std::numeric_limits<index_t>::max(),
   matxDropDim     = cuda::std::numeric_limits<index_t>::max() - 1,
   matxEnd         = cuda::std::numeric_limits<index_t>::max() - 2,

diff --git a/include/matx/core/error.h b/include/matx/core/error.h
@@ -117,16 +117,23 @@ namespace matx
     default:
       return "Unknown";
     };
-
-    return "Unknown";
   }
 
   namespace detail {
   struct matxException : public std::exception
   {
     matxError_t e;
     char str[400];
-    std::stringstream stack;
+    // Stack trace captured as std::string rather than std::stringstream.
+    //
+    // std::stringstream deletes its copy constructor, making any class that
+    // holds one non-copyable. MSVC introduced warning C5272 (promoted to an
+    // error by /WX) in VS 18 (MSVC 14.4x / 2026) that fires when a
+    // non-copyable type is thrown: the C++ standard requires a thrown object
+    // to be copyable because the runtime may need to copy it during stack
+    // unwinding. std::string is copyable, avoids the warning entirely, and is
+    // semantically equivalent here since stack is only ever read as a string.
+    std::string stack;
 
     /**
      * @brief Throw an exception and print a stack trace
@@ -141,15 +148,19 @@ namespace matx
     {
       snprintf(str, sizeof(str), "matxException (%s: %s) - %s:%d\n",
                matxErrorString(error), s, file, line);
-      detail::printStackTrace(stack);
+      std::ostringstream oss;
+      detail::printStackTrace(oss);
+      stack = oss.str();
     }
 
     matxException(matxError_t error, const std::string &s, const char *file, int line)
         : e(error)
     {
       snprintf(str, sizeof(str), "matxException (%s: %s) - %s:%d\n",
                matxErrorString(error), s.c_str(), file, line);
-      detail::printStackTrace(stack);
+      std::ostringstream oss;
+      detail::printStackTrace(oss);
+      stack = oss.str();
     }
 
     const char* what() const noexcept override { return str; }
@@ -182,7 +193,7 @@ namespace matx
   catch (matx::detail::matxException & e)                       \
   {                                                             \
     MATX_LOG_FATAL("{}", e.what());                             \
-    MATX_LOG_FATAL("Stack Trace:\n{}", e.stack.str());          \
+    MATX_LOG_FATAL("Stack Trace:\n{}", e.stack);               \
     exit(1);                                                    \
   }