CUDA: сумма всех элементов в массиве с использованием линеаризованной 2D совместно используемой памяти

Я новичок в CUDA, и алгоритмы в целом. Может кто-нибудь сказать мне, правильно ли я это делаю или есть ли лучший способ сделать это. Одной из проблем является то, что ввод и вывод кода должны осуществляться на графическом процессоре, чтобы не было копирования памяти между хостом и устройством.

#include "cuda_runtime.h"#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdint.h>

#include <iostream>

#define TILE_WIDTH 8

__global__ void gpu_sumElements(int height, int width, float *in, float *out){

extern __shared__ float cache[];

int w = blockIdx.x * blockDim.x + threadIdx.x; // Col // width
int h = blockIdx.y * blockDim.y + threadIdx.y;

int index = h * width + w;
int cacheIndex = threadIdx.y * blockDim.x + threadIdx.x;

float temp = 0;

if ((w < width) && (h < height)){
temp += in[index];
//index += (height * width);
}

cache[cacheIndex] = temp;
__syncthreads();

int i = (blockDim.x * blockDim.y) / 2;
while (i != 0){
if (cacheIndex < i)
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
i /= 2;
}

if (cacheIndex == 0)
out[blockIdx.y * gridDim.x + blockIdx.x] = cache[0];
}int main(){
// Initial Parameters
int width = 2363;
int height = 781;
float my_sum = 0;

int block_height = (height - 1) / TILE_WIDTH + 1;
int block_width = (width - 1) / TILE_WIDTH + 1;

dim3 dimGrid(block_width, block_height, 1);
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);

int sharedMemSize = TILE_WIDTH * TILE_WIDTH * sizeof(float);
// Initialize host arrays
float *test_array = new float[height * width];
float *out = new float[height * width];
for (int i = 0; i < (height * width); i++)
test_array[i] = 1.0f;
// Initialize device arrays
float *gpu_temp_array;
float *gpu_out;
cudaMalloc((void **)&gpu_temp_array, (height * width * sizeof(float)));
cudaMalloc((void **)&gpu_out, (height * width * sizeof(float)));
cudaMemcpy(gpu_out, test_array, (height * width * sizeof(float)), cudaMemcpyHostToDevice);

// Copy these, need them elsewhere
float sum_height = height;
float sum_width = width ;
dim3 sum_dimGrid = dimGrid;
int i = (height * width);

// Launch kernel, get & print results
while (i != 0){
gpu_sumElements<<<sum_dimGrid, dimBlock, sharedMemSize>>>(sum_height, sum_width, gpu_out, gpu_temp_array);
cudaMemcpy(gpu_out, gpu_temp_array, (sum_height * sum_width * sizeof(float)), cudaMemcpyDeviceToDevice);
cudaMemset(gpu_temp_array, 0, (height * width * sizeof(float)));

sum_height = ceil(sum_height/TILE_WIDTH);
sum_width = ceil(sum_width/TILE_WIDTH);;
sum_dimGrid.x = (sum_width - 1) / TILE_WIDTH + 1;
sum_dimGrid.y = (sum_height - 1) / TILE_WIDTH + 1;
i /= TILE_WIDTH*TILE_WIDTH;
}

cudaMemcpy(out, gpu_out, (height * width * sizeof(float)), cudaMemcpyDeviceToHost);
std::cout << out[0] << std::endl << std::endl;

delete[] test_array;
delete[] out;
cudaFree(gpu_out);
cudaFree(gpu_temp_array);

system("pause");
return 0;
}