Адаптируйте код к сложению целочисленных матриц. Найдите наилучшую конфигурацию выполнения. Объясните.Заранее спасибо

#include <cuda_runtime.h>
#include <stdio.h>
#define WIN32_LEAN_AND_MEAN
#include <Windows.h>
#include <stdint.h>
#define CHECK(call) 
{ 
const cudaError_t error = call; 
if (error != cudaSuccess) 
{ 
printf("Ошибка: %s:%d, ", __FILE__, __LINE__); 
printf("код:%d, причина: %s\n", error, cudaGetErrorString(error)); 
exit(-10*error); 
} 
} 
typedef struct timeval { long tv_sec; long tv_usec;
} timeval;
int gettimeofday(struct timeval* tp, struct timezone* tzp)
{
 // Это магическое число - количество 100-наносекундных интервалов с 1 января // 1601 (UTC) до 00:00:00 1 января 1970 static const uint64_t EPOCH = ((uint64_t)116444736000000000ULL);
 SYSTEMTIME system_time; FILETIME file_time; uint64_t time;
 GetSystemTime(&system_time); SystemTimeToFileTime(&system_time, &file_time); time = ((uint64_t)file_time.dwLowDateTime); time += ((uint64_t)file_time.dwHighDateTime) << 32;
 tp->tv_sec = (long)((time - EPOCH) / 10000000L); tp->tv_usec = (long)(system_time.wMilliseconds * 1000); return 0;
}
double cpuSecond() { struct timeval tp; gettimeofday(&tp, NULL); return ((double)tp.tv_sec + (double)tp.tv_usec * 1.e-6);
}
void initialData(float* ip, const int size)
{
 int i;
 for (i = 0; i < size; i++) { ip[i] = (float)(rand() & 0xFF) / 10.0f; }
 return;
}
void sumMatrixOnHost(float* A, float* B, float* C, const int nx, const int ny)
{
 float* ia = A; float* ib = B; float* ic = C;
 for (int iy = 0; iy < ny; iy++) {
 for (int ix = 0; ix < nx; ix++) { ic[ix] = ia[ix] + ib[ix];
 }
 ia += nx; ib += nx; ic += nx; }
 return;
}
void checkResult(float* hostRef, float* gpuRef, const int N)
{
 double epsilon = 1.0E-8; bool match = 1;
 for (int i = 0; i < N; i++) { if (abs(hostRef[i] - gpuRef[i]) > epsilon) { match = 0; printf("host %f gpu %f\n", hostRef[i], gpuRef[i]); break; } }
 if (match) printf("Arrays match.\n\n"); else printf("Arrays do not match.\n\n");
}
// 2D сетка и 2D блок
__global__ void sumMatrixOnGPU2D(float* MatA, float* MatB, float* MatC, int nx, int ny)
{
 unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x; unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y; unsigned int idx = iy * nx + ix;
 if (ix < nx && iy < ny) MatC[idx] = MatA[idx] + MatB[idx];
}
int main(int argc, char** argv)
{
 printf("%s Starting...\n", argv[0]);
 // настраиваем устройство int dev = 0; cudaDeviceProp deviceProp; CHECK(cudaGetDeviceProperties(&deviceProp, dev)); printf("Using Device %d: %s\n", dev, deviceProp.name); CHECK(cudaSetDevice(dev));
 // устанавливаем объём данных матрицы int nx = 1 << 14; int ny = 1 << 14;
 int nxy = nx * ny; int nBytes = nxy * sizeof(float); printf("Matrix size: nx %d ny %d\n", nx, ny);
 // выделяем память хоста float* h_A, * h_B, * hostRef, * gpuRef; h_A = (float*)malloc(nBytes); h_B = (float*)malloc(nBytes); hostRef = (float*)malloc(nBytes); gpuRef = (float*)malloc(nBytes);
 // инициализируем данные на стороне хоста double iStart = cpuSecond(); initialData(h_A, nxy); initialData(h_B, nxy); double iElaps = cpuSecond() - iStart; printf("Matrix initialization elapsed %f sec\n", iElaps);
 memset(hostRef, 0, nBytes); memset(gpuRef, 0, nBytes);
 // складываем матрицы на хосте для проверки результатов iStart = cpuSecond(); sumMatrixOnHost(h_A, h_B, hostRef, nx, ny); iElaps = cpuSecond() - iStart; printf("sumMatrixOnHost elapsed %f sec\n", iElaps);
 // выделяем глобальную память устройства float* d_MatA, * d_MatB, * d_MatC; CHECK(cudaMalloc((void**)&d_MatA, nBytes)); CHECK(cudaMalloc((void**)&d_MatB, nBytes)); CHECK(cudaMalloc((void**)&d_MatC, nBytes));
 // передаём данные с хоста на устройство CHECK(cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice)); CHECK(cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice));
 // вызываем ядро на стороне хоста int dimx = 32; int dimy = 32; dim3 block(dimx, dimy); dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);
 iStart = cpuSecond(); sumMatrixOnGPU2D << <grid, block >> > (d_MatA, d_MatB, d_MatC, nx, ny); CHECK(cudaDeviceSynchronize()); iElaps = cpuSecond() - iStart; printf("sumMatrixOnGPU2D <<<(%d,%d), (%d,%d)>>> elapsed %f sec\n", grid.x, grid.y, block.x, block.y, iElaps); // проверяем ошибку ядра CHECK(cudaGetLastError());
 // копируем результаты на хост CHECK(cudaMemcpy(gpuRef, d_MatC, nBytes, cudaMemcpyDeviceToHost));
 // проверяем результаты устройства checkResult(hostRef, gpuRef, nxy);
 // очищаем глобальную память устройства CHECK(cudaFree(d_MatA)); CHECK(cudaFree(d_MatB)); CHECK(cudaFree(d_MatC));
 // освобождаем память хоста free(h_A); free(h_B); free(hostRef); free(gpuRef);
 // сбрасываем устройство CHECK(cudaDeviceReset());
 return (0);
}

#include <stdlib.h>
#include "cuda_runtime.h"
#include <stddef.h>
#include <stdarg.h>
#include <stdint.h>
#include <stdatomic.h>
#include <stdio.h>
#include <stdint.h>
#include <math.h>
#include <cuBLAS.h>
#include <cuRAND.h>
#include <iostream>
#include <device_launch_parameters.h>
#define WIN32_LEAN_AND_MEAN
#include <Windows.h>
#define CHECK(call)
{ 
    const cudaError_t error = call; 
    if (error != cudaSuccess) 
    { 
        printf("Ошибка: %s:%d, ", __FILE__, __LINE__); 
        printf("код:%d, причина: %s\n", error, cudaGetErrorString(error)); 
        exit(-10*error); 
    } 
} 
typedef struct timeval {
    long tv_sec; long tv_usec;
} timeval;
 
int gettimeofday(struct timeval* tp, struct timezone* tzp)
{
    // Это магическое число - количество 100-наносекундных интервалов с 1 января 
    // 1601 (UTC) до 00:00:00 1 января 1970 
    static const uint64_t EPOCH = ((uint64_t)116444736000000000ULL);
 
    SYSTEMTIME system_time; 
    FILETIME file_time;
    uint64_t time;
 
    GetSystemTime(&system_time);
    SystemTimeToFileTime(&system_time, &file_time);
    time = ((uint64_t)file_time.dwLowDateTime);
    time += ((uint64_t)file_time.dwHighDateTime) << 32;
 
    tp->tv_sec = (long)((time - EPOCH) / 10000000L);
    tp->tv_usec = (long)(system_time.wMilliseconds * 1000); 
    return 0;
}
double cpuSecond() {
    struct timeval tp;
 gettimeofday(&tp, NULL);
    return ((double)tp.tv_sec + (double)tp.tv_usec * 1.e-6);
}
void initialData(float* ip, const int size)
{
    int i;
    for (i = 0; i < size; i++) 
    { 
        ip[i] = (float)(rand() & 0xFF) / 10.0f;
    }
    return;
}
void sumMatrixOnHost(float* A, float* B, float* C, const int nx, const int ny)
{
    float* ia = A;
    float* ib = B; 
    float* ic = C;
    for (int iy = 0; iy < ny; iy++) 
    {
        for (int ix = 0; ix < nx; ix++) 
        {
            ic[ix] = ia[ix] + ib[ix];
        }
        ia += nx;
        ib += nx;
        ic += nx;
    }
    return;
}
void checkResult(float* hostRef, float* gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;
    for (int i = 0; i < N; i++)
    { 
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        { 
            match = 0;
            printf("host %f gpu %f\n", hostRef[i], gpuRef[i]);
            break;
        }
    }
    if (match) 
        printf("Arrays match.\n\n");
    else printf("Arrays do not match.\n\n");
}
 
// 2D сетка и 2D блок
__global__ void sumMatrixOnGPU2D(float* MatA, float* MatB, float* MatC, int nx, int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;
    if (ix < nx && iy < ny)
        MatC[idx] = MatA[idx] + MatB[idx];
}
int main(int argc, char** argv)
{
    printf("%s Starting...\n", argv[0]);
    
    // настраиваем устройство
     int dev = 0;
     cudaDeviceProp deviceProp;
     CHECK(cudaGetDeviceProperties(&deviceProp, dev));
     printf("Using Device %d: %s\n", dev, deviceProp.name);
     CHECK(cudaSetDevice(dev));
    
     // устанавливаем объём данных матрицы
     int nx = 1 << 14;
     int ny = 1 << 14;
 
    int nxy = nx * ny;
    int nBytes = nxy * sizeof(float);
    printf("Matrix size: nx %d ny %d\n", nx, ny);
    
    // выделяем память хоста
     float* h_A, * h_B, * hostRef, * gpuRef;
     h_A = (float*)malloc(nBytes);
     h_B = (float*)malloc(nBytes);
     hostRef = (float*)malloc(nBytes);
     gpuRef = (float*)malloc(nBytes);
    
     // инициализируем данные на стороне хоста 
     double iStart = cpuSecond();
     initialData(h_A, nxy);
     initialData(h_B, nxy);
     double iElaps = cpuSecond() - iStart;
     printf("Matrix initialization elapsed %f sec\n", iElaps);
 
     memset(hostRef, 0, nBytes);
     memset(gpuRef, 0, nBytes);
    
    // складываем матрицы на хосте для проверки результатов 
    iStart = cpuSecond();
    sumMatrixOnHost(h_A, h_B, hostRef, nx, ny);
    iElaps = cpuSecond() - iStart;
    printf("sumMatrixOnHost elapsed %f sec\n", iElaps);
 
    
    // выделяем глобальную память устройства
     float* d_MatA, * d_MatB, * d_MatC;
     CHECK(cudaMalloc((void**)&d_MatA, nBytes));
     CHECK(cudaMalloc((void**)&d_MatB, nBytes));
     CHECK(cudaMalloc((void**)&d_MatC, nBytes));
    
     // передаём данные с хоста на устройство
     CHECK(cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice));
     CHECK(cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice));
    
     // вызываем ядро на стороне хоста 
     int dimx = 32;
     int dimy = 32;
     dim3 block(dimx, dimy);
     dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);
    iStart = cpuSecond();
    sumMatrixOnGPU2D << <grid, block >> > (d_MatA, d_MatB, d_MatC, nx, ny);
    CHECK(cudaDeviceSynchronize());
    iElaps = cpuSecond() - iStart;
    printf("sumMatrixOnGPU2D <<<(%d, %d), (%d, %d)>>> elapsed %f sec\n", grid.x, grid.y, block.x, block.y, iElaps);
    
    // проверяем ошибку ядра
     CHECK(cudaGetLastError());
 
    // копируем результаты на хост
    CHECK(cudaMemcpy(gpuRef, d_MatC, nBytes, cudaMemcpyDeviceToHost));
    // проверяем результаты устройства
     checkResult(hostRef, gpuRef, nxy);
    // очищаем глобальную память устройства
     CHECK(cudaFree(d_MatA));
     CHECK(cudaFree(d_MatB));
     CHECK(cudaFree(d_MatC));
 
    // освобождаем память хоста
     free(h_A);
     free(h_B);
     free(hostRef);
     free(gpuRef);
 
    // сбрасываем устройство
     CHECK(cudaDeviceReset());
    return (0);
}