hey I have this kernel which runs perfectly fine in release mode but doesn't in debug mode, can you take a look at it. The kernel itself is fine, the way I'm calling it with is probably not fine.#define CEIL_DIV(m,n) ( (m) + (n) - 1 ) / (n)
#define FETCH_FLOAT4(pointer) (reinterpret_cast<float4 *>(&(pointer))[0])


template< int TLIE_M_PER_BLOCK,
    int TLIE_N_PER_BLOCK,
    int TLIE_K_PER_BLOCK,
    int TLIE_M_PER_THREAD,
    int TLIE_N_PER_THREAD >
__global__ void im2winConvCUDAimplentNCHWHPCkernel(float* inptr, float* filptr, float* outptr, size_t* d_dims_a, size_t* d_dims_b, size_t* d_dims_c) {
    const int THREAD_M_PER_BLOCK = TLIE_M_PER_BLOCK / TLIE_M_PER_THREAD;
    const int THREAD_N_PER_BLOCK = TLIE_N_PER_BLOCK / TLIE_N_PER_THREAD;
    const int THREAD_PER_BLOCK = THREAD_M_PER_BLOCK * THREAD_N_PER_BLOCK;

    const int K = d_dims_b[1] * d_dims_b[2] * d_dims_b[3];

    const int tid = threadIdx.x;
    const int tx = tid / THREAD_N_PER_BLOCK;
    const int ty = tid % THREAD_N_PER_BLOCK;
    const int bx = blockIdx.x;
    const int by = blockIdx.y;

    const int input_batch = d_dims_a[0];
    const int input_channel = d_dims_a[1];
    const int input_height = d_dims_a[2];
    const int input_width = d_dims_a[3];

    const int filter_batch = d_dims_b[0];
    const int filter_channel = d_dims_b[1];
    const int filter_height = d_dims_b[2];
    const int filter_width = d_dims_b[3];

    const int output_batch = d_dims_c[0];
    const int output_channel = d_dims_c[1];
    const int output_height = d_dims_c[2];
    const int output_width = d_dims_c[3];

    const int stride_height = (input_height - filter_height) / (output_height - 1);
    const int stride_width = (input_width - filter_width) / (output_width - 1);

    __shared__ float Asmem[2][TLIE_K_PER_BLOCK * TLIE_N_PER_BLOCK];
    __shared__ float Bsmem[2][TLIE_K_PER_BLOCK * TLIE_M_PER_BLOCK];

    float Areg[2][TLIE_N_PER_THREAD];
    float Breg[2][TLIE_M_PER_THREAD];
    float Creg[TLIE_M_PER_THREAD * TLIE_N_PER_THREAD];
    memset(Creg, 0, sizeof(Creg));

    const int THREAD_TILE_A_PER_COL = TLIE_K_PER_BLOCK;
    const int THREAD_TILE_B_PER_COL = TLIE_K_PER_BLOCK;

    const int TILE_A_ROW = tid % THREAD_TILE_A_PER_COL;
    const int TILE_B_ROW = tid % THREAD_TILE_B_PER_COL;

    const int TILE_A_COL = tid / THREAD_TILE_A_PER_COL;
    const int TILE_B_COL = tid / THREAD_TILE_B_PER_COL;

    const int TILE_A_COL_STRIDE = THREAD_PER_BLOCK / THREAD_TILE_A_PER_COL;
    const int TILE_B_COL_STRIDE = THREAD_PER_BLOCK / THREAD_TILE_B_PER_COL;

    if (THREAD_PER_BLOCK % THREAD_TILE_A_PER_COL != 0 || THREAD_PER_BLOCK % THREAD_TILE_B_PER_COL != 0) { printf("TILE_THREAD_STRIDE ERROR"); return; }
    //if (threadIdx.x == 0 && threadIdx.y == 0)
    //{
    //    printf("first thread of each block\n");
    //}
#pragma unroll
    for (int i = 0; i < TLIE_M_PER_BLOCK; i += TILE_B_COL_STRIDE) {
        int id_oc = bx * TLIE_M_PER_BLOCK + TILE_B_COL + i;
        int id_ic = (0 + TILE_B_ROW) / (filter_height * filter_width);
        int res_ic = (0 + TILE_B_ROW) % (filter_height * filter_width);
        int id_fh = res_ic % filter_height;
        int id_fw = res_ic / filter_height;

        Bsmem[0][(TILE_B_ROW)*TLIE_M_PER_BLOCK + TILE_B_COL + i] =
            filptr[id_oc * filter_channel * filter_height * filter_width
            + id_ic * filter_height * filter_width
            + id_fh
            + id_fw * filter_height];
    }

#pragma unroll
    for (int j = 0; j < TLIE_N_PER_BLOCK; j += TILE_A_COL_STRIDE) {
        int id_ob = (by * TLIE_N_PER_BLOCK + TILE_A_COL + j) / (output_height * output_width);
        int res_ob = (by * TLIE_N_PER_BLOCK + TILE_A_COL + j) % (output_height * output_width);
        int id_oh = res_ob / output_width;
        int id_ow = res_ob % output_width;

        int id_ic = (0 + TILE_A_ROW) / (filter_height * filter_width);
        int res_ic = (0 + TILE_A_ROW) % (filter_height * filter_width);
        int id_fh = res_ic % filter_height;
        int id_fw = res_ic / filter_height;

        Asmem[0][(TILE_A_ROW)*TLIE_N_PER_BLOCK + TILE_A_COL + j] =
            inptr[id_ob * input_channel * output_height * filter_height * input_width
            + id_ic * output_height * filter_height * input_width
            + id_oh * filter_height * input_width
            + id_ow * stride_width * filter_height
            + id_fh
            + id_fw * filter_height];
        if (threadIdx.x == 0 && blockIdx.x == 0 && blockIdx.y == 0)
        {
            printf("output[%d] \n", id_ob * input_channel * output_height * filter_height * input_width
                + id_ic * output_height * filter_height * input_width
                + id_oh * filter_height * input_width
                + id_ow * stride_width * filter_height
                + id_fh
                + id_fw * filter_height);
        }
    }
    __syncthreads();

#pragma unroll
    for (int i = 0; i < TLIE_M_PER_THREAD; i += 4) {
        FETCH_FLOAT4(Breg[0][i]) = FETCH_FLOAT4(Bsmem[0][0 * TLIE_M_PER_BLOCK + tx * TLIE_M_PER_THREAD + i]);
    }
#pragma unroll
    for (int j = 0; j < TLIE_N_PER_THREAD; j += 4) {
        FETCH_FLOAT4(Areg[0][j]) = FETCH_FLOAT4(Asmem[0][0 * TLIE_N_PER_BLOCK + ty * TLIE_N_PER_THREAD + j]);
    }

    int write_stage_idx = 1;
    int k_count = 0;
    do {
        k_count += TLIE_K_PER_BLOCK;
        int load_stage_idx = write_stage_idx ^ 1;

#pragma unroll
        for (int inner_k = 0; inner_k < TLIE_K_PER_BLOCK - 1; ++inner_k) {
#pragma unroll
            for (int i = 0; i < TLIE_M_PER_THREAD; i += 4) {
                FETCH_FLOAT4(Breg[(inner_k + 1) % 2][i]) = FETCH_FLOAT4(Bsmem[load_stage_idx][(inner_k + 1) * TLIE_M_PER_BLOCK + tx * TLIE_M_PER_THREAD + i]);
            }
#pragma unroll
            for (int j = 0; j < TLIE_N_PER_THREAD; j += 4) {
                FETCH_FLOAT4(Areg[(inner_k + 1) % 2][j]) = FETCH_FLOAT4(Asmem[load_stage_idx][(inner_k + 1) * TLIE_N_PER_BLOCK + ty * TLIE_N_PER_THREAD + j]);
            }
#pragma unroll
            for (int i = 0; i < TLIE_M_PER_THREAD; ++i) {
#pragma unroll
                for (int j = 0; j < TLIE_N_PER_THREAD; ++j) {
                    Creg[i * TLIE_N_PER_THREAD + j] +=
                        Breg[inner_k % 2][i] * Areg[inner_k % 2][j];
                }
            }
        }

        if (k_count < K) {
#pragma unroll
            for (int i = 0; i < TLIE_M_PER_BLOCK; i += TILE_B_COL_STRIDE) {
                int id_oc = bx * TLIE_M_PER_BLOCK + TILE_B_COL + i;
                int id_ic = (k_count + TILE_B_ROW) / (filter_height * filter_width);
                int res_ic = (k_count + TILE_B_ROW) % (filter_height * filter_width);
                int id_fh = res_ic % filter_height;
                int id_fw = res_ic / filter_height;

                Bsmem[write_stage_idx][(TILE_B_ROW)*TLIE_M_PER_BLOCK + TILE_B_COL + i] =
                    filptr[id_oc * filter_channel * filter_height * filter_width
                    + id_ic * filter_height * filter_width
                    + id_fh
                    + id_fw * filter_height];
            }

#pragma unroll
            for (int j = 0; j < TLIE_N_PER_BLOCK; j += TILE_A_COL_STRIDE) {
                int id_ob = (by * TLIE_N_PER_BLOCK + TILE_A_COL + j) / (output_height * output_width);
                int res_ob = (by * TLIE_N_PER_BLOCK + TILE_A_COL + j) % (output_height * output_width);
                int id_oh = res_ob / output_width;
                int id_ow = res_ob % output_width;

                int id_ic = (k_count + TILE_A_ROW) / (filter_height * filter_width);
                int res_ic = (k_count + TILE_A_ROW) % (filter_height * filter_width);
                int id_fh = res_ic % filter_height;
                int id_fw = res_ic / filter_height;

                Asmem[write_stage_idx][(TILE_A_ROW)*TLIE_N_PER_BLOCK + TILE_A_COL + j] =
                    inptr[id_ob * input_channel * output_height * filter_height * input_width
                    + id_ic * output_height * filter_height * input_width
                    + id_oh * filter_height * input_width
                    + id_ow * stride_width * filter_height
                    + id_fh
                    + id_fw * filter_height];
            }
            __syncthreads();
            write_stage_idx ^= 1;
        }

#pragma unroll
        for (int i = 0; i < TLIE_M_PER_THREAD; i += 4) {
            FETCH_FLOAT4(Breg[0][i]) = FETCH_FLOAT4(Bsmem[load_stage_idx ^ 1][0 + tx * TLIE_M_PER_THREAD + i]);
        }
#pragma unroll
        for (int j = 0; j < TLIE_N_PER_THREAD; j += 4) {
            FETCH_FLOAT4(Areg[0][j]) = FETCH_FLOAT4(Asmem[load_stage_idx ^ 1][0 + ty * TLIE_N_PER_THREAD + j]);
        }

#pragma unroll
        for (int i = 0; i < TLIE_M_PER_THREAD; ++i) {
#pragma unroll
            for (int j = 0; j < TLIE_N_PER_THREAD; ++j) {
                Creg[i * TLIE_N_PER_THREAD + j] +=
                    Breg[1][i] * Areg[1][j];

                //if (Creg[i * TLIE_N_PER_THREAD + j] - 0.022827f <= 0.0001)
                //{
                //    printf("index : %d \n", i* TLIE_N_PER_THREAD + j);
                //}
            }
        }
    } while (k_count < K);

#pragma unroll
    for (int i = 0; i < TLIE_M_PER_THREAD; ++i) {
#pragma unroll
        for (int j = 0; j < TLIE_N_PER_THREAD; ++j) {
            int id_oc = bx * TLIE_M_PER_BLOCK + tx * TLIE_M_PER_THREAD + i;
            int id_ob = (by * TLIE_N_PER_BLOCK + ty * TLIE_N_PER_THREAD + j) / (output_height * output_width);
            int res_ob = (by * TLIE_N_PER_BLOCK + ty * TLIE_N_PER_THREAD + j) % (output_height * output_width);
            int id_oh = res_ob / output_width;
            int id_ow = res_ob % output_width;
            //if (threadIdx.x == 0 && blockIdx.x == 0 && blockIdx.y == 0)
            //{
            //    printf("output[%d] \n", id_ob * output_channel * output_height * output_width
            //        + id_oc * output_height * output_width
            //        + id_oh * output_width
            //        + id_ow);
            //}
            outptr[id_ob * output_channel * output_height * output_width
                + id_oc * output_height * output_width
                + id_oh * output_width
                + id_ow] = Creg[i * TLIE_N_PER_THREAD + j];
        }
    }
    //printf("Idx_x : %d\n", Idx_x);
    // printf("Idx_y : %d\n", Idx_y);
}

int im2winConvCUDAimplentNCHWHPC(float* inptr, float* filptr, float* outptr, size_t* dims_a, size_t* dims_b, size_t* dims_c) {
    //set memory
    // float *d_inptr, *d_filptr, *d_outptr;
    size_t* d_dims_a, * d_dims_b, * d_dims_c;
    //alloca memory
    // cudaMalloc(&d_inptr, sizeof(float)*dims_a[0]*dims_a[1]*dims_a[2]*dims_a[3]);
    // cudaMalloc(&d_filptr, sizeof(float)*dims_b[0]*dims_b[1]*dims_b[2]*dims_b[3]);
    // cudaMalloc(&d_outptr, sizeof(float)*dims_c[0]*dims_c[1]*dims_c[2]*dims_c[3]);
    cudaMalloc(&d_dims_a, sizeof(size_t) * 4);
    cudaMalloc(&d_dims_b, sizeof(size_t) * 4);
    cudaMalloc(&d_dims_c, sizeof(size_t) * 4);
    //init input and filter
    // cudaMemcpy(d_inptr, inptr, sizeof(float)*dims_a[0]*dims_a[1]*dims_a[2]*dims_a[3], cudaMemcpyHostToDevice);
    // cudaMemcpy(d_filptr, filptr, sizeof(float)*dims_b[0]*dims_b[1]*dims_b[2]*dims_b[3], cudaMemcpyHostToDevice);
    cudaMemcpy(d_dims_a, dims_a, sizeof(size_t) * 4, cudaMemcpyHostToDevice);
    cudaMemcpy(d_dims_b, dims_b, sizeof(size_t) * 4, cudaMemcpyHostToDevice);
    cudaMemcpy(d_dims_c, dims_c, sizeof(size_t) * 4, cudaMemcpyHostToDevice);

    //set blockSize and gridSize
    int output_batch = dims_c[0];
    int output_channel = dims_c[1];
    int output_height = dims_c[2];
    int output_width = dims_c[3];

    int filter_batch = dims_b[0];
    int filter_channel = dims_b[1];
    int filter_height = dims_b[2];
    int filter_width = dims_b[3];
    if (filter_channel < 8) filter_channel = 8;

    const int M = output_channel;
    const int N = output_batch * output_height * output_width;
    const int K = filter_channel * filter_height * filter_width;
    constexpr int TLIE_M_PER_BLOCK = 32;
    constexpr int TLIE_N_PER_BLOCK = 32;
    constexpr int TLIE_K_PER_BLOCK = 8;
    constexpr int TLIE_M_PER_THREAD = 8;
    constexpr int TLIE_N_PER_THREAD = 8;
    const dim3 blockSize((TLIE_M_PER_BLOCK / TLIE_M_PER_THREAD) * (TLIE_N_PER_BLOCK / TLIE_N_PER_THREAD));
    const dim3 gridSize(CEIL_DIV(M, TLIE_M_PER_BLOCK),
        CEIL_DIV(N, TLIE_N_PER_BLOCK));

    //const dim3 blockSize(16, 1, 1);
    //const dim3 gridSize(1, 90, 1);

    im2winConvCUDAimplentNCHWHPCkernel< TLIE_M_PER_BLOCK,
        TLIE_N_PER_BLOCK,
        TLIE_K_PER_BLOCK,
        TLIE_M_PER_THREAD,
        TLIE_N_PER_THREAD > << <gridSize, blockSize >> >
        (inptr, filptr, outptr, d_dims_a, d_dims_b, d_dims_c);
    //result to host
    //cudaMemcpy(outptr, d_outptr, sizeof(float)*dims_c[0]*dims_c[1]*dims_c[2]*dims_c[3], cudaMemcpyDeviceToHost);
    //clean device memory
    cudaFree(d_dims_a);
    cudaFree(d_dims_b);
    cudaFree(d_dims_c);
    return 0;
}int main() {

    const int channels = 1;
    const int height = 32;
    const int width = 32;

    float* new_input;

    // test for the rearrange function
    const int x1 = (height - 3)/1 + 1;
    const int x2 = (width - 3)/1 + 1;
    const int d = x1 * x2;
    
    float* temp_input = new float[d];
    float* new_temp;

    for (size_t i = 0; i < d; i++)
    {
        temp_input[i] = static_cast<float>(i);
    }

    //printf("testing data ------------------------ \n");
    //for (size_t i = 0; i < x1; i++)
    //{
    //    for (size_t j = 0; j < x2; j++)
    //    {
    //        printf("%f ", temp_input[j + i * x1]);
    //    }
    //    printf("\n");
    //}
    //printf("after rearrange ------------------------ \n");
    //new_temp = rearrangeChannels(temp_input, 1, x1, x2, 3, 1, 1, 0);
    //int t1 = (x1 + 2 * 0 - 3) / 1 + 1;

    //for (size_t i = 0; i < t1; i++)
    //{
    //    for (size_t j = 0; j < 3 * x2; j++)
    //    {
    //        printf("%f ", new_temp[j + i * 3 * x2]);
    //    }
    //    printf("\n");
    //}

    float input[] = {
        0.0f, 1.0f, 2.0f,
        10.0f, 11.0f, 12.0f,
        20.0f, 21.0f, 22.0f,
        100.0f, 101.0f, 102.0f,
        110.0f, 111.0f, 112.0f,
        120.0f, 121.0f, 122.0f,
    };

    


    // Define dimensions
    size_t dims_a[4] = { 1, 1, height, width };  // Input: 1 batch, 3 channels, 32x32 image
    size_t dims_a_n[4] = { 1, 1, x1, 3 * width };  // Input: 1 batch, 3 channels, 32x32 image
    size_t dims_b[4] = { 1, 1, 3, 3 };    // Filter: 1 batch, 3 channels, 3x3 kernel
    size_t dims_c[4] = { 1, 1, x1, x2 };  // Output: 1 batch, 1 channel, 30x30 image

    int ho = (dims_a[2] + 2 * 0 - dims_b[3]) / 1 + 1;
    int wo = (dims_a[3] + 2 * 0 - dims_b[3]) / 1 + 1;
    int windows = ho * dims_a[3]/ 3;
    int remains = ho * dims_a[3] - windows * 3;
    int new_height = ho;
    int new_width = width * 3;

    // Allocate host memory
    std::vector<float> h_input(dims_a[0] * dims_a[1] * dims_a[2] * dims_a[3]);
    std::vector<float> h_filter(dims_b[0] * dims_b[1] * dims_b[2] * dims_b[3]);
    std::vector<float> h_output_cpu(dims_c[0] * dims_c[1] * dims_c[2] * dims_c[3]);
    std::vector<float> h_output_gpu(dims_c[0] * dims_c[1] * dims_c[2] * dims_c[3]);

    // Initialize input and filter with random data
    std::random_device rd;
    std::mt19937 gen(rd());
    std::uniform_real_distribution<> dis(-1.0, 1.0);
    for (auto& val : h_input) val = dis(gen);
    for (auto& val : h_filter) val = dis(gen);

    for (size_t i = 0; i < h_input.size(); i++)
    {
        h_input[i] = (static_cast<float>(i)/h_input.size());
        
    }
    for (size_t i = 0; i < h_filter.size(); i++)
    {
        h_filter[i] = (static_cast<float>(i) / h_filter.size());
    }
    
    //for (size_t i = 0; i < dims_a[2]; i++)
    //{
    //    for (size_t j = 0; j < dims_a[3]; j++)
    //    {
    //        printf("%f ", h_input[i * dims_a[3] + j]);
    //    }
    //    printf("\n");
    //}

    new_input = rearrangeChannels(h_input.data(), dims_a[1], dims_a[2], dims_a[3], 3, 1, 1, 0);

    // Allocate device memory
    float* d_input, * d_filter, * d_output;
    size_t* d_dims_a, * d_dims_b, * d_dims_c, *d_dims_a_n;

    cudaMalloc(&d_input, new_height * new_width * sizeof(float));
    //cudaMalloc(&d_input, h_input.size() * sizeof(float));
    cudaMalloc(&d_filter, h_filter.size() * sizeof(float));
    cudaMalloc(&d_output, h_output_gpu.size() * sizeof(float));
    cudaMalloc(&d_dims_a, 4 * sizeof(size_t));
    cudaMalloc(&d_dims_a_n, 4 * sizeof(size_t));
    cudaMalloc(&d_dims_b, 4 * sizeof(size_t));
    cudaMalloc(&d_dims_c, 4 * sizeof(size_t));

    // Copy data to device
    cudaMemcpy(d_input, new_input, new_height * new_width * sizeof(float), cudaMemcpyHostToDevice);
    //cudaMemcpy(d_input, h_input.data(), h_input.size() * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_filter, h_filter.data(), h_filter.size() * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_dims_a, dims_a, 4 * sizeof(size_t), cudaMemcpyHostToDevice);
    cudaMemcpy(d_dims_a_n, dims_a_n, 4 * sizeof(size_t), cudaMemcpyHostToDevice);
    cudaMemcpy(d_dims_b, dims_b, 4 * sizeof(size_t), cudaMemcpyHostToDevice);
    cudaMemcpy(d_dims_c, dims_c, 4 * sizeof(size_t), cudaMemcpyHostToDevice);

    // Launch kernel
    const int blocksize = 15;
    dim3 block(blocksize * blocksize);
    dim3 grid(
        (dims_c[2] + blocksize - 1) / blocksize,
        (dims_c[3] + blocksize - 1) / blocksize,
        dims_c[0]
    );

    //dim3 grid((dims_c[2] + block.x - 1) / block.x, (dims_c[3] + block.x - 1) / block.x, dims_c[0]);
    //im2winKernelConvNCHWBASE<blocksize> << <grid, block >> > (d_input, d_filter, d_output, d_dims_a, d_dims_b, d_dims_c);
    //im2winConvCUDAimplentNCHWBASE(h_input.data(), h_filter.data(), h_output_gpu.data(), dims_a, dims_b, dims_c);
    //im2winKernelConvNCHWBASE<blocksize> << <grid, block >> > (d_input, d_filter, d_output, d_dims_a_n, d_dims_b, d_dims_c);
    //  WORKS : -> down
    // im2winKernelConvNCHWBASE<blocksize> << <grid, block >> > (d_input, d_filter, d_output, d_dims_a, d_dims_b, d_dims_c);
    im2winConvCUDAimplentNCHWHPC(d_input, d_filter, d_output, dims_a, dims_b, dims_c);
    // Copy result back to host
    cudaMemcpy(h_output_gpu.data(), d_output, h_output_gpu.size() * sizeof(float), cudaMemcpyDeviceToHost);

    // Run CPU implementation
    cpuImplementation(new_input, h_filter.data(), h_output_cpu.data(), dims_a, dims_b, dims_c);

    // Compare results
    float max_diff = 0.0f;
    for (size_t i = 0; i < h_output_cpu.size(); ++i) {
        float diff = std::abs(h_output_cpu[i] - h_output_gpu[i]);
        if (diff > max_diff) max_diff = diff;
        if (diff > 0.001)
        {
            printf("output index : %d\n", i);
            break;
        }
    }
    printf("out : %f\n", h_output_gpu[1]);
    printf("out : %f\n", h_output_cpu[1]);
    std::cout << "Max difference between CPU and GPU results: " << max_diff << std::endl;

    // Clean up
    cudaFree(d_input);
    cudaFree(d_filter);
    cudaFree(d_output);
    cudaFree(d_dims_a);
    cudaFree(d_dims_b);
    cudaFree(d_dims_c);

    return 0;
} data has to be rearranged so that is fine.