CUDA implements matrix replication

  The key is thread allocation , A matrix that divides a large matrix into pieces , Threaded ny,nx To represent the rows and columns of the matrix respectively , But because the general matrix is stored in a one-dimensional array , So finally, it corresponds to the global memory index You can also convert rows and columns into one-dimensional index.

#include <stdio.h>

const int N = 128;  // matrix' dim
const int TILE_DIM = 32; // block size
const int SIZE = sizeof(int) * N * N; // memory

void __global__ copy(int *da, int *db, const int N);

void __global__ copy(int *da, int *db, const int N)
{
    
    int ny = blockIdx.y * TILE_DIM + threadIdx.y;
    int nx = blockIdx.x * TILE_DIM + threadIdx.x;

    int index = ny*N + nx; // get each index
    if(ny<N && nx<N)
    {
    
        db[index] = da[index];
    }


}

int main(int argc, char *argv[])
{
    
    // host memory and assignment
    int *ha, *hb;
    ha = (int *)malloc(SIZE);
    hb = (int *)malloc(SIZE);

    for(int i=0; i< N*N; ++i)
    {
    
        ha[i] = 100;
    }
    // device memry
    int *da, *db;
    cudaMalloc((void **)&da, SIZE);
    cudaMalloc((void **)&db, SIZE);

    cudaMemcpy(da, ha, SIZE, cudaMemcpyHostToDevice);
    // kernel function
    const dim3 block_size(TILE_DIM, TILE_DIM);
    const int grid_size_x = (N +TILE_DIM -1) / TILE_DIM;
    const int grid_size_y =  grid_size_x;
    const dim3 grid_size(grid_size_x, grid_size_y);

    copy<<<grid_size,block_size>>>(da,db,N);

    // device to host
    cudaMemcpy(hb,db,SIZE,cudaMemcpyDeviceToHost);

    printf("%d\n",hb[100]);

    // free
    free(ha);
    free(hb);
    cudaFree(da);
    cudaFree(db);

    return 0;

}