Hey, here in this code, cuda type float4 doesn't have an operator = overload, can you fix this so that the errors go away: #define OFFSET(row, col, ld) ((row) * (ld) + (col))

// transfer float4
//#define FETCH_FLOAT4(pointer) (reinterpret_cast<float4*>(&(pointer))[0])
//#define FETCH_FLOAT4(pointer) (reinterpret_cast<float4*>(const_cast<void*>(&pointer))[0])
#define FETCH_FLOAT4(pointer) (*reinterpret_cast<const float4*>(&pointer))

template <
    const int BLOCK_SIZE_M,  // width of block of C that each thread block calculate
    const int BLOCK_SIZE_K,  // height of block of A that each thread block load into shared memory
    const int BLOCK_SIZE_N,  // height of block of C that each thread block calculate
    const int THREAD_SIZE_Y, // height of block of C that each thread calculate
    const int THREAD_SIZE_X,  // width of block of C that each thread calculate
    const bool ENABLE_DOUBLE_BUFFER // whether enable double buffering or not
>
__global__ void FusedIm2ColDoubleBufferMatMul(
    const float* __restrict__ input,
    const float* __restrict__ kernel,
    float* __restrict__ output,
    const int C, const int H, const int W,
    const int K, const int P, const int S,
    const int OH, const int OW) {

    // Block index
    int bx = blockIdx.x;
    int by = blockIdx.y;

    // Thread index
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    // size of thread block
    const int bszx = BLOCK_SIZE_N / THREAD_SIZE_X;
    const int bszy = BLOCK_SIZE_M / THREAD_SIZE_Y;
    const int THREAD_NUM_PER_BLOCK = bszy * bszx;

    // thread id
    const int tid = ty * bszx + tx;

    // shared memory
    __shared__ float As[BLOCK_SIZE_M * 2][BLOCK_SIZE_K];
    __shared__ float Bs[BLOCK_SIZE_K * 2][BLOCK_SIZE_N];

    // registers for C
    float accum[THREAD_SIZE_Y][THREAD_SIZE_X] = { 0 };
    // registers for A and B
    float frag_a[THREAD_SIZE_Y];
    float frag_b[THREAD_SIZE_X];

    // threads needed to load one row of tile
    // / 4 is because float4 is used
    const int A_TILE_THREAD_PER_ROW = BLOCK_SIZE_K / 4;
    const int B_TILE_THREAD_PER_ROW = BLOCK_SIZE_N / 4;

    // row number and col number that needs to be loaded by this thread
    const int A_TILE_ROW_START = tid / A_TILE_THREAD_PER_ROW;
    const int B_TILE_ROW_START = tid / B_TILE_THREAD_PER_ROW;

    const int A_TILE_COL = tid % A_TILE_THREAD_PER_ROW * 4;
    const int B_TILE_COL = tid % B_TILE_THREAD_PER_ROW * 4;

    // row stride that thread uses to load multiple rows of a tile
    const int A_TILE_ROW_STRIDE = THREAD_NUM_PER_BLOCK / A_TILE_THREAD_PER_ROW;
    const int B_TILE_ROW_STRIDE = THREAD_NUM_PER_BLOCK / B_TILE_THREAD_PER_ROW;

    int A_SM_OFFSET = 0;
    int B_SM_OFFSET = 0;

    const int tile_size = C * P * P;

    for (int tile_idx = 0; tile_idx < tile_size; tile_idx += BLOCK_SIZE_K) {
        // load A (input) from global memory to shared memory
#pragma unroll
        for (int i = 0; i < BLOCK_SIZE_M; i += A_TILE_ROW_STRIDE) {
            int row = BLOCK_SIZE_M * by + A_TILE_ROW_START + i;
            int col = A_TILE_COL + tile_idx;
            if (row < K && col < tile_size) {
                int c = col / (P * P);
                int pix = col % (P * P);
                int kh = pix / P;
                int kw = pix % P;
                int oh = row / OW;
                int ow = row % OW;
                int h = oh * S + kh;
                int w = ow * S + kw;
                if (h >= 0 && h < H && w >= 0 && w < W) {
                    FETCH_FLOAT4(As[A_TILE_ROW_START + i + A_SM_OFFSET][A_TILE_COL]) = FETCH_FLOAT4(input[OFFSET(c, h * W + w, H * W)]);
                }
                else {
                    FETCH_FLOAT4(As[A_TILE_ROW_START + i + A_SM_OFFSET][A_TILE_COL]) = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
                }
            }
        }

        // load B (kernel) from global memory to shared memory
#pragma unroll
        for (int i = 0; i < BLOCK_SIZE_K; i += B_TILE_ROW_STRIDE) {
            int row = tile_idx + B_TILE_ROW_START + i;
            int col = B_TILE_COL + BLOCK_SIZE_N * bx;
            if (row < tile_size && col < K) {
                FETCH_FLOAT4(Bs[B_TILE_ROW_START + i + B_SM_OFFSET][B_TILE_COL]) = FETCH_FLOAT4(kernel[OFFSET(col, row, tile_size)]);
            }
            else {
                FETCH_FLOAT4(Bs[B_TILE_ROW_START + i + B_SM_OFFSET][B_TILE_COL]) = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
            }
        }

        __syncthreads();

        // compute c
#pragma unroll
        for (int k = 0; k < BLOCK_SIZE_K; ++k) {
            // load A from shared memory to register
#pragma unroll
            for (int thread_y = 0; thread_y < THREAD_SIZE_Y; ++thread_y) {
                frag_a[thread_y] = As[ty * THREAD_SIZE_Y + thread_y + A_SM_OFFSET][k];
            }

            // load B from shared memory to register
#pragma unroll
            for (int thread_x = 0; thread_x < THREAD_SIZE_X; thread_x += 4) {
                FETCH_FLOAT4(frag_b[thread_x]) = FETCH_FLOAT4(Bs[k + B_SM_OFFSET][THREAD_SIZE_X * tx + thread_x]);
            }

#pragma unroll
            for (int thread_y = 0; thread_y < THREAD_SIZE_Y; ++thread_y) {
#pragma unroll
                for (int thread_x = 0; thread_x < THREAD_SIZE_X; ++thread_x) {
                    accum[thread_y][thread_x] += frag_a[thread_y] * frag_b[thread_x];
                }
            }
        }
        // change offset so new block of A and B will not override the previous data
        // therefore __syncthreads() is not needed;
        A_SM_OFFSET ^= BLOCK_SIZE_M;
        B_SM_OFFSET ^= BLOCK_SIZE_K;
    }

    // store back to C (output)
#pragma unroll
    for (int thread_y = 0; thread_y < THREAD_SIZE_Y; ++thread_y) {
#pragma unroll
        for (int thread_x = 0; thread_x < THREAD_SIZE_X; ++thread_x) {
            int row = BLOCK_SIZE_M * by + ty * THREAD_SIZE_Y + thread_y;
            int col = BLOCK_SIZE_N * bx + tx * THREAD_SIZE_X + thread_x;
            if (row < K * OH * OW && col < K) {
                output[OFFSET(col, row, K * OH * OW)] = accum[thread_y][thread_x];
            }
        }
    }
}
The lines where stuff like this FETCH_FLOAT4(frag_b[thread_x]) = FETCH_FLOAT4(Bs[k + B_SM_OFFSET][THREAD_SIZE_X * tx + thread_x]); will fail because don't have an operator for float4 = float4