having trouble running this code, with this error : Severity	Code	Description	Project	File	Line	Suppression State
Error	MSB3721	The command ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.8\bin\nvcc.exe" -gencode=arch=compute_52,code=\"sm_52,compute_52\" --use-local-env -ccbin "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Tools\MSVC\14.33.31629\bin\HostX64\x64" -x cu   -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.8\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.8\include"     --keep-dir x64\Release  -maxrregcount=0  --machine 64 --compile -cudart static    -DWIN32 -DWIN64 -DNDEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /O2 /Fdx64\Release\vc143.pdb /FS   /MD " -o A:\DEV\CUDA\TrajectoryGenerator\TrajectoryGenerator\x64\Release\optimized.cu.obj "A:\DEV\CUDA\TrajectoryGenerator\TrajectoryGenerator\optimized.cu"" exited with code 255.	TrajectoryGenerator	C:\Program Files\Microsoft Visual Studio\2022\Enterprise\MSBuild\Microsoft\VC\v170\BuildCustomizations\CUDA 11.8.targets	785	 #define SMEM_LDA (36)  // Adjusted for 32x32 (32 + padding)
#define SMEM_LDB (32)

// Helper functions (unchanged)
__device__ __forceinline__ void ldg32_nc_0(float& reg, const void* ptr) {
    asm volatile("{.reg .pred p;\\n"
        "mov.b32 %0, 0;\\n"
        "ld.global.nc.f32 %0, [%1];}\\n"
        : "=f"(reg)
        : "l"(ptr));
}

__device__ __forceinline__ uint32_t smem_u32addr(const void* smem_ptr) {
    uint32_t addr;
    asm("{.reg .u64 u64addr;\\n"
        " cvta.to.shared.u64 u64addr, %1;\\n"
        " cvt.u32.u64 %0;}\\n"
        : "=r"(addr)
        : "l"(smem_ptr));
    return addr;
}

__device__ __forceinline__ void lds128(float& reg0, float& reg1, float& reg2, float& reg3, const uint32_t& addr) {
    asm volatile("ld.shared.v4.f32 {%0, %1, %2, %3}, [%4];\\n"
        : "=f"(reg0), "=f"(reg1), "=f"(reg2), "=f"(reg3)
        : "r"(addr));
}

__device__ __forceinline__ void sts128(const float& reg0, const float& reg1, const float& reg2, const float& reg3, const uint32_t& addr) {
    asm volatile("st.shared.v4.f32 [%0], {%1, %2, %3, %4};\\n"
        :
    : "r"(addr), "f"(reg0), "f"(reg1), "f"(reg2), "f"(reg3));
}

__device__ __forceinline__ void sts32(const float& reg, const uint32_t& addr) {
    asm volatile("st.shared.f32 [%0], %1;\\n" : : "r"(addr), "f"(reg));
}

// Kernel function
__global__ void sgemm_32x32x32(int m, int n, int k,
    const float* a,
    const float* b,
    float* c,
    const float* bias) {
    __shared__ __align__(16 * 1024) char smem[6 * 1024]; // Reduced shared memory

    float* ashare = reinterpret_cast<float*>(smem);
    float* bshare = reinterpret_cast<float*>(smem + 4 * 1024); // 4k shared mem for A, 2k for B
    float sum[4][4] = { 0 };  // Reduced from 8x8 to 4x4
    float panelA[2][4] = { 0 }, panelB[2][4] = { 0 };  // Reduced from 8 to 4

    int from_a = (blockIdx.y * 32 + threadIdx.x / 8 * 4) * k + threadIdx.x % 8;
    int from_b = (threadIdx.x / 32) * n + blockIdx.x * 32 + threadIdx.x % 32;

    float a_ldg_reg[4], b_ldg_reg[4];

    uint32_t a_sts_addr = smem_u32addr(ashare + (threadIdx.x % 8) * SMEM_LDA + (threadIdx.x / 8) * 4);
    uint32_t b_sts_addr = smem_u32addr(bshare + (threadIdx.x / 32) * SMEM_LDB + (threadIdx.x % 32));

    uint32_t aptr_base = smem_u32addr(ashare + (threadIdx.x / 16) * 4);
    uint32_t bptr_base = smem_u32addr(bshare + (threadIdx.x % 16) * 4);

    // Load first
    if (blockIdx.y * 32 + threadIdx.x / 8 * 4 < m && threadIdx.x % 8 < k) {
        for (int i = 0; i < 4 && blockIdx.y * 32 + threadIdx.x / 8 * 4 + i < m; ++i) {
            ldg32_nc_0(a_ldg_reg[i], (const char*)(a + from_a) + i * k * sizeof(float));
        }
        sts128(a_ldg_reg[0], a_ldg_reg[1], a_ldg_reg[2], a_ldg_reg[3], a_sts_addr);
    }

    if (threadIdx.x / 32 < k && blockIdx.x * 32 + threadIdx.x % 32 < n) {
        for (int i = 0; i < 4 && threadIdx.x / 32 + i < k; ++i) {
            ldg32_nc_0(b_ldg_reg[i], (const char*)(b + from_b) + i * 32 * sizeof(float));
        }
        for (int i = 0; i < 4 && threadIdx.x / 32 + i < k; ++i) {
            sts32(b_ldg_reg[i], b_sts_addr + i * 32 * sizeof(float));
        }
    }
    __syncthreads();

    from_a += 8;
    from_b += 8 * n;
    a_sts_addr ^= 0x1000;  // Adjusted for smaller shared memory
    b_sts_addr ^= 0x800;   // Adjusted for smaller shared memory

    // Load first panel
    lds128(panelA[0][0], panelA[0][1], panelA[0][2], panelA[0][3], aptr_base);
    lds128(panelB[0][0], panelB[0][1], panelB[0][2], panelB[0][3], bptr_base);

    for (int loop = 0; loop < k; loop += 8) {
#pragma unroll
        for (int subk = 0; subk < 8; ++subk) {
            if (7 == subk && loop < k - 8) {
                if (blockIdx.y * 32 + threadIdx.x / 8 * 4 < m && threadIdx.x % 8 + loop + 8 < k) {
                    sts128(a_ldg_reg[0], a_ldg_reg[1], a_ldg_reg[2], a_ldg_reg[3], a_sts_addr);
                }
                if (threadIdx.x / 32 + loop + 8 < k && blockIdx.x * 32 + threadIdx.x % 32 < n) {
                    for (int i = 0; i < 4 && threadIdx.x / 32 + loop + 8 + i < k; ++i) {
                        sts32(b_ldg_reg[i], b_sts_addr + i * 32 * sizeof(float));
                    }
                }
                __syncthreads();
                from_a += 8;
                from_b += 8 * n;
                aptr_base ^= 0x1000;  // Adjusted for smaller shared memory
                bptr_base ^= 0x800;   // Adjusted for smaller shared memory
                a_sts_addr ^= 0x1000; // Adjusted for smaller shared memory
                b_sts_addr ^= 0x800;  // Adjusted for smaller shared memory
            }

            const int pp = (subk + 1) % 2;
            lds128(panelA[pp][0], panelA[pp][1], panelA[pp][2], panelA[pp][3],
                aptr_base + ((subk + 1) % 8) * SMEM_LDA * sizeof(float));
            lds128(panelB[pp][0], panelB[pp][1], panelB[pp][2], panelB[pp][3],
                bptr_base + ((subk + 1) % 8) * SMEM_LDB * sizeof(float));

            if (0 == subk && loop < k - 8) {
                if (blockIdx.y * 32 + threadIdx.x / 8 * 4 < m && threadIdx.x % 8 + loop + 8 < k) {
                    for (int i = 0; i < 4 && blockIdx.y * 32 + threadIdx.x / 8 * 4 + i < m; ++i) {
                        ldg32_nc_0(a_ldg_reg[i], (const char*)(a + from_a) + i * k * sizeof(float));
                    }
                }
                if (threadIdx.x / 32 + loop + 8 < k && blockIdx.x * 32 + threadIdx.x % 32 < n) {
                    for (int i = 0; i < 4 && threadIdx.x / 32 + loop + 8 + i < k; ++i) {
                        ldg32_nc_0(b_ldg_reg[i], (const char*)(b + from_b) + i * 32 * sizeof(float));
                    }
                }
            }

#pragma unroll
            for (int i = 0; i < 4; ++i) {
#pragma unroll
                for (int j = 0; j < 4; ++j) {
                    sum[i][j] += panelA[subk % 2][i] * panelB[subk % 2][j];
                }
            }
        }
    }

    int write_offset = (blockIdx.y * 32 + (threadIdx.x / 16) * 4) * n + blockIdx.x * 32 + (threadIdx.x % 16) * 4;
    if (blockIdx.y * 32 + (threadIdx.x / 16) * 4 < m && blockIdx.x * 32 + (threadIdx.x % 16) * 4 < n) {
        for (int i = 0; i < 4 && blockIdx.y * 32 + (threadIdx.x / 16) * 4 + i < m; ++i) {
            for (int j = 0; j < 4 && blockIdx.x * 32 + (threadIdx.x % 16) * 4 + j < n; ++j) {
                c[write_offset + i * n + j] = sum[i][j] + bias[blockIdx.x * 32 + (threadIdx.x % 16) * 4 + j];
            }
        }
    }
}

// Function to launch the kernel
void tiledMatMulBias32x32x32(float* A, float* B, float* C, float* bias, int M, int N, int K) {
    dim3 grid((N + 31) / 32, (M + 31) / 32);
    sgemm_32x32x32 << <grid, 256 >> > (M, N, K, A, B, C, bias);
}

// cuBLAS implementation with built-in bias addition
void cublasMatMulBias(cublasHandle_t handle, float* A, float* B, float* C, float* bias, int M, int N, int K) {
    const float alpha = 1.0f;
    const float beta = 1.0f;

    // Perform matrix multiplication with bias addition
    cublasSgemm(handle,
        CUBLAS_OP_N, CUBLAS_OP_N,
        N, M, K,
        &alpha,
        B, N,
        A, K,
        &beta,
        C, N);

    // Add bias to the result
    cublasSaxpy(handle, M * N, &alpha, bias, 1, C, 1);
}

// CPU implementation for ground truth
void cpuMatMulBias(const std::vector<float>& A, const std::vector<float>& B, std::vector<float>& C,
    const std::vector<float>& bias, int M, int N, int K) {
    for (int i = 0; i < M; ++i) {
        for (int j = 0; j < N; ++j) {
            float sum = 0.0f;
            for (int k = 0; k < K; ++k) {
                sum += A[i * K + k] * B[k * N + j];
            }
            C[i * N + j] = sum + bias[j];
        }
    }
}

// Function to initialize matrix with random values
void initializeMatrix(std::vector<float>& matrix, int size) {
    std::random_device rd;
    std::mt19937 gen(rd());
    std::uniform_real_distribution<> dis(-1.0, 1.0);
    for (int i = 0; i < size; ++i) {
        matrix[i] = dis(gen);
    }
}

// Function to compare results
bool compareResults(const std::vector<float>& a, const std::vector<float>& b, float tolerance = 1e-5) {
    if (a.size() != b.size()) return false;
    for (size_t i = 0; i < a.size(); ++i) {
        if (std::abs(a[i] - b[i]) > tolerance) return false;
    }
    return true;
}

int main() {
    // Define matrix dimensions
    int M1 = 12, K1 = 3, N1 = 32;
    int M2 = 12, K2 = 32, N2 = 64;
    int M3 = 12, K3 = 64, N3 = 128;

    // Initialize matrices and biases
    std::vector<float> A1(M1 * K1), B1(K1 * N1), C1(M1 * N1), bias1(N1);
    std::vector<float> B2(K2 * N2), C2(M2 * N2), bias2(N2);
    std::vector<float> B3(K3 * N3), C3(M3 * N3), bias3(N3);

    initializeMatrix(A1, A1.size());
    initializeMatrix(B1, B1.size());
    initializeMatrix(B2, B2.size());
    initializeMatrix(B3, B3.size());
    initializeMatrix(bias1, bias1.size());
    initializeMatrix(bias2, bias2.size());
    initializeMatrix(bias3, bias3.size());

    // Allocate device memory
    float* d_A1, * d_B1, * d_C1, * d_bias1;
    float* d_B2, * d_C2, * d_bias2;
    float* d_B3, * d_C3, * d_bias3;

    cudaMalloc(&d_A1, A1.size() * sizeof(float));
    cudaMalloc(&d_B1, B1.size() * sizeof(float));
    cudaMalloc(&d_C1, C1.size() * sizeof(float));
    cudaMalloc(&d_bias1, bias1.size() * sizeof(float));
    cudaMalloc(&d_B2, B2.size() * sizeof(float));
    cudaMalloc(&d_C2, C2.size() * sizeof(float));
    cudaMalloc(&d_bias2, bias2.size() * sizeof(float));
    cudaMalloc(&d_B3, B3.size() * sizeof(float));
    cudaMalloc(&d_C3, C3.size() * sizeof(float));
    cudaMalloc(&d_bias3, bias3.size() * sizeof(float));

    // Copy data to device
    cudaMemcpy(d_A1, A1.data(), A1.size() * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B1, B1.data(), B1.size() * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_bias1, bias1.data(), bias1.size() * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B2, B2.data(), B2.size() * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_bias2, bias2.data(), bias2.size() * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B3, B3.data(), B3.size() * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_bias3, bias3.data(), bias3.size() * sizeof(float), cudaMemcpyHostToDevice);

    // Custom CUDA implementation (32x32x32 tiled)
    auto start = std::chrono::high_resolution_clock::now();

    tiledMatMulBias32x32x32(d_A1, d_B1, d_C1, d_bias1, M1, N1, K1);
    tiledMatMulBias32x32x32(d_C1, d_B2, d_C2, d_bias2, M2, N2, K2);
    tiledMatMulBias32x32x32(d_C2, d_B3, d_C3, d_bias3, M3, N3, K3);

    cudaDeviceSynchronize();
    auto end = std::chrono::high_resolution_clock::now();
    auto duration = std::chrono::duration_cast<std::chrono::microseconds>(end - start);
    std::cout << "Tiled 32x32x32 CUDA time: " << duration.count() << " microseconds" << std::endl;

    // Copy result back to host
    std::vector<float> tiled_result(C3.size());
    cudaMemcpy(tiled_result.data(), d_C3, C3.size() * sizeof(float), cudaMemcpyDeviceToHost);

    // cuBLAS implementation
    cublasHandle_t handle;
    cublasCreate(&handle);

    start = std::chrono::high_resolution_clock::now();

    cublasMatMulBias(handle, d_A1, d_B1, d_C1, d_bias1, M1, N1, K1);
    cublasMatMulBias(handle, d_C1, d_B2, d_C2, d_bias2, M2, N2, K2);
    cublasMatMulBias(handle, d_C2, d_B3, d_C3, d_bias3, M3, N3, K3);

    cudaDeviceSynchronize();
    end = std::chrono::high_resolution_clock::now();
    duration = std::chrono::duration_cast<std::chrono::microseconds>(end - start);
    std::cout << "cuBLAS time: " << duration.count() << " microseconds" << std::endl;

    // Copy result back to host
    std::vector<float> cublas_result(C3.size());
    cudaMemcpy(cublas_result.data(), d_C3, C3.size() * sizeof(float), cudaMemcpyDeviceToHost);

    cublasDestroy(handle);

    // CPU implementation for ground truth
    std::vector<float> cpu_C1(M1 * N1), cpu_C2(M2 * N2), cpu_C3(M3 * N3);

    start = std::chrono::high_resolution_clock::now();

    cpuMatMulBias(A1, B1, cpu_C1, bias1, M1, N1, K1);
    cpuMatMulBias(cpu_C1, B2, cpu_C2, bias2, M2, N2, K2);
    cpuMatMulBias(cpu_C2, B3, cpu_C3, bias3, M3, N3, K3);

    end = std::chrono::high_resolution_clock::now();
    duration = std::chrono::duration_cast<std::chrono::microseconds>(end - start);
    std::cout << "CPU time: " << duration.count() << " microseconds" << std::endl;

    // Compare results
    bool tiled_matches = compareResults(tiled_result, cpu_C3);
    bool cublas_matches = compareResults(cublas_result, cpu_C3);

    std::cout << "Tiled 32x32x32 CUDA result matches CPU: " << (tiled_matches ? "Yes" : "No") << std::endl;
    std::cout << "cuBLAS result matches CPU: " << (cublas_matches ? "Yes" : "No") << std::endl;

    // Free device memory
    cudaFree(d_A1);
    cudaFree(d_B1);
    cudaFree(d_C1);
    cudaFree(d_bias1);
    cudaFree(d_B2);
    cudaFree(d_C2);
    cudaFree(d_bias2);
    cudaFree(d_B3);
    cudaFree(d_C3);
    cudaFree(d_bias3);

    return 0;
}