Fixing Windows compile with standard C++

[loscrossos]

I propose a change to standard C++ code that also fixes the windows build.
I also subtmitted a PR with said changes.

There are 2 issues i would like to adress for this to happen: both break compilation in windows for awhile already but AFAICS both are easily adressed by adhering to C++ standards:

"and" operators in v1.2.2.post1

This release introduces "and" operators (instead of '&&') that MSVC does not recognize in causal-conv1d\csrc\causal_conv1d.cpp
This is fixed by including the cISO646 standard header, that allows MSVC to recognize them.
I would think that in order to use them you have to use that header anyway since its part of the standard:

https://en.cppreference.com/w/cpp/header/ciso646

but alas i put it in a windows guard.

csrc\causal_conv1d.cpp

#ifdef _WIN32
#include <ciso646>  // Enables 'and', 'or', 'not' keywords on MSVC
#endif

preprocessor directive in a macro

introduced in v1.3.0.post1

Starting with this version there are 2 cu files that do not compile for me.
This happened when adding support for ROCm using a if-directive inside of a macro.

The error occurs because the BOOL_SWITCH macro is expanding into a lambda that contains preprocessor directives for ROCm/CUDA differentiation (ifndef, else, endif), but PDs cannot appear inside a lambda function's body in C++. directives must appear at file/function scope.

At first i was sure that a preprocessing directive must be at the beginning of a line. There is even a MSVC error that says so:

https://learn.microsoft.com/en-us/cpp/error-messages/compiler-errors-1/compiler-error-c2014?view=msvc-170

But that is not the case. The current standard says otherwise:

https://gcc.gnu.org/onlinedocs/cpp.pdf

Page 7 chapter 1.4 The preprocessing language

Yet still on the same page:

"The '#' which begins a directive cannot come from a macro expansion."

And this is where it breaks. As understand it, you cannot place ifdef, endif, etc. inside macro expansions. The compiler aborts before it can even complain that the directive has whitespaces in front of it (not sure if it would dare to). I would think the code should have not compiled at all but it seems it did for a couple of releases already.

I changed the code by copying the code into 2 blocks simply moving the if-directives outside the macro. not my proudest code since the only difference is a single "(void *)" but it keeps it simple and you can see with the naked eye that nothing breaks.

Here the changed blocks:

csrc\causal_conv1d_fwd.cu

template<int kNThreads, int kWidth, typename input_t, typename weight_t>
void causal_conv1d_fwd_launch(ConvParamsBase &params, cudaStream_t stream) {
    static constexpr int kNElts = sizeof(input_t) == 4 ? 4 : 8;
    
    #ifndef USE_ROCM
    BOOL_SWITCH(params.seqlen % kNElts == 0, kIsVecLoad, [&] {
        using Ktraits = Causal_conv1d_fwd_kernel_traits<kNThreads, kWidth, kIsVecLoad, input_t, weight_t>;
        constexpr int kSmemSize = Ktraits::kSmemSize;
        dim3 grid(params.batch, params.dim);
        auto kernel = &causal_conv1d_fwd_kernel<Ktraits>;
        if (kSmemSize >= 48 * 1024) {
            C10_CUDA_CHECK(cudaFuncSetAttribute(
                kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, kSmemSize));
        }
        kernel<<<grid, Ktraits::kNThreads, kSmemSize, stream>>>(params);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    });
    #else
    BOOL_SWITCH(params.seqlen % kNElts == 0, kIsVecLoad, [&] {
        using Ktraits = Causal_conv1d_fwd_kernel_traits<kNThreads, kWidth, kIsVecLoad, input_t, weight_t>;
        constexpr int kSmemSize = Ktraits::kSmemSize;
        dim3 grid(params.batch, params.dim);
        auto kernel = &causal_conv1d_fwd_kernel<Ktraits>;
        if (kSmemSize >= 48 * 1024) {
            // There is a slight signature discrepancy in HIP and CUDA "FuncSetAttribute" function.
            C10_CUDA_CHECK(cudaFuncSetAttribute(
                (void *) kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, kSmemSize));
            std::cerr << "Warning (causal_conv1d fwd launch): attempting to set maxDynamicSharedMemorySize on an AMD GPU which is currently a non-op (in ROCm versions <= 6.1). This might lead to undefined behavior. \n" << std::endl;
        }
        kernel<<<grid, Ktraits::kNThreads, kSmemSize, stream>>>(params);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    });
    #endif
}

csrc\causal_conv1d_bwd.cu

template<int kNThreads, int kWidth, typename input_t, typename weight_t>
void causal_conv1d_bwd_launch(ConvParamsBwd &params, cudaStream_t stream) {
    static constexpr int kNElts = sizeof(input_t) == 4 ? 4 : 8;
    #ifndef USE_ROCM
    BOOL_SWITCH(params.seqlen % kNElts == 0, kIsVecLoad, [&] {
        BOOL_SWITCH(params.silu_activation, kSiluAct, [&] {
            using Ktraits = Causal_conv1d_bwd_kernel_traits<kNThreads, kWidth, kSiluAct, kIsVecLoad, input_t, weight_t>;
            constexpr int kSmemSize = Ktraits::kSmemSize;
            dim3 grid(params.batch, params.dim);
            auto kernel = &causal_conv1d_bwd_kernel<Ktraits>;
            if (kSmemSize >= 48 * 1024) {
                C10_CUDA_CHECK(cudaFuncSetAttribute(
                    kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, kSmemSize));
            }
            kernel<<<grid, Ktraits::kNThreads, kSmemSize, stream>>>(params);
            C10_CUDA_KERNEL_LAUNCH_CHECK();
        });
    });
    #else
    BOOL_SWITCH(params.seqlen % kNElts == 0, kIsVecLoad, [&] {
        BOOL_SWITCH(params.silu_activation, kSiluAct, [&] {
            using Ktraits = Causal_conv1d_bwd_kernel_traits<kNThreads, kWidth, kSiluAct, kIsVecLoad, input_t, weight_t>;
            constexpr int kSmemSize = Ktraits::kSmemSize;
            dim3 grid(params.batch, params.dim);
            auto kernel = &causal_conv1d_bwd_kernel<Ktraits>;

            if (kSmemSize >= 48 * 1024) {
                // There is a slight signature discrepancy in HIP and CUDA "FuncSetAttribute" function.
                C10_CUDA_CHECK(cudaFuncSetAttribute(
                    (void *) kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, kSmemSize));
                std::cerr << "Warning (causal_conv1d bwd launch): attempting to set maxDynamicSharedMemorySize on an AMD GPU which is currently a non-op (in ROCm versions <= 6.1). This might lead to undefined behavior. \n" << std::endl;
            }
            kernel<<<grid, Ktraits::kNThreads, kSmemSize, stream>>>(params);
            C10_CUDA_KERNEL_LAUNCH_CHECK();
        });
    });
    #endif
}

i pushed this in #58