How to use __shared__ in CUDA?

My code is suppsed to work like this:

-take in_martix of NxN elements and R factor

-it should give back a matrix of size [N-2R]*[N-2R] with each element being a sum of in_matrix elements in R radius it should work like this for N=4 R=1

Now I am trying to understand how to use shared memory to speed things up and I want to clarify some things.

#include "cuda_runtime.h"
#include <stdio.h>
#include <iostream>
#include <cuda_profiler_api.h>

#define N 8
#define R 2
#define K 1
#define THREAD_BLOCK_SIZE 32
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char* file, int line, bool abort = true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}
using namespace std;

__global__ void MatrixStencil(int* d_tab_begin, int* d_out_begin, int d_N, int d_R, int d_K) {
    int tx = threadIdx.x + blockIdx.x * blockDim.x;
    int ty = threadIdx.y + blockIdx.y * blockDim.y;
    
    int out_local = 0;
    
    __shared__ int tabS[N*N];
    tabS[tx * N + ty] = *(d_tab_begin + ty * d_N + tx);
    __syncthreads();
        if ((tx  < d_N - 2 * d_R) && (ty  < d_N - 2 * d_R)) {
            for (int col = tx; col <= tx + 2 * d_R; col++)
                for (int row = ty; row <= ty + 2 * d_R; row++)
                    out_local += tabS[col * N + row];
                    //out_local += *(d_tab_begin + col * d_N + row);
                    

            *(d_out_begin + (tx) * (d_N - 2 * R) + ty) = out_local;
        }
    

}
void random_ints(int tab[N*N]) {
    for (int i = 0; i < N; i++)
        for (int j = 0; j < N; j++)
            tab[i*N + j] = 1;
}

int main() {

    static int tab[N*N];
    random_ints(tab);

    int tab_size = sizeof(int) * N * N;
    int out_size = sizeof(int) * (N - 2 * R) * (N - 2 * R);

    dim3 BLOCK_SIZE(THREAD_BLOCK_SIZE, THREAD_BLOCK_SIZE);
    dim3 GRID_SIZE(ceil((float)N / (float)(THREAD_BLOCK_SIZE  )), ceil((float)N / (float)(THREAD_BLOCK_SIZE  )));

    void** d_tab;
    void** d_out;

    cudaMalloc((void**)&d_tab, tab_size);
    cudaMalloc((void**)&d_out, out_size);

    cudaMemcpyAsync(d_tab, tab, tab_size, cudaMemcpyHostToDevice,0);

    int* d_tab_begin = (int*)(d_tab);
    int* d_out_begin = (int*)(d_out);

    MatrixStencil << < GRID_SIZE, BLOCK_SIZE>> > (d_tab_begin, d_out_begin, N, R, K);
    gpuErrchk(cudaPeekAtLastError());
    

    int* out = (int*)malloc(out_size);
    
    cudaMemcpyAsync(out, d_out, out_size, cudaMemcpyDeviceToHost,0);
    cudaDeviceSynchronize();

    for (int col = 0; col < N - 2 * R; col++)
    {
        for (int row = 0; row < N - 2 * R; row++)
        {
            cout << *(out + ((col * (N - 2 * R)) + row)) << " ";
        }
        cout << endl;
    }
    
}

Does it work like this? : each thread loads one element of matrix to shared matrix, after that we want to synchronize threads so that we are sure every element is loaded. In our case tabS should be loaded with 1's.

    tabS[tx * N + ty] = *(d_tab_begin + ty * d_N + tx);
    __syncthreads();

And here is the part that I don't get:

out_local += tabS[col * N + row];
//out_local += *(d_tab_begin + col * d_N + row);

Shoudln't both versions work properly? Apparenlty only the commented one works. The shared one gives back some random numbers or 0's. (do I do some illegal memory accesses?)