From b0a88a57c14ccb0b3eba6e1fb9de861aeb360cbf Mon Sep 17 00:00:00 2001 From: lucasdelimanogueira Date: Thu, 2 May 2024 14:08:28 -0300 Subject: [PATCH 1/2] Fix sum tensor reduce cuda v1 naive --- norch/csrc/cuda.cu | 70 +++++++++++++++++++++++++++++++++------------- 1 file changed, 50 insertions(+), 20 deletions(-) diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index c757724..fb0ade6 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -5,7 +5,6 @@ #define THREADS_PER_BLOCK 128 #define TILE_SIZE 32 -#define SHMEM_SIZE THREADS_PER_BLOCK * sizeof(float) __host__ void cpu_to_cuda(Tensor* tensor) { @@ -60,41 +59,72 @@ __host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_da } -__global__ void sum_tensor_cuda_kernel(float* data, float* result_data) { +__global__ void sum_tensor_cuda_kernel(float* data, float* result_data, int size) { + __shared__ float partial_sum[THREADS_PER_BLOCK]; - __shared__ int partial_sum[SHMEM_SIZE]; + int tid = threadIdx.x; + int i = blockIdx.x * blockDim.x + threadIdx.x; - // each thread loads one element from global to shared mem - // note use of 1D thread indices (only) in this kernel - int i = blockIdx.x*blockDim.x + threadIdx.x; - - partial_sum[threadIdx.x] = data[i]; + partial_sum[tid] = (i < size) ? data[i] : 0; __syncthreads(); - // do reduction in shared mem - for (int s=1; s < blockDim.x; s *=2) - { - if (threadIdx.x % (2 * s) == 0) { - partial_sum[threadIdx.x] += partial_sum[threadIdx.x + s]; + // Perform block-wise reduction + for (int s = blockDim.x / 2; s > 0; s >>= 1) { + if (tid < s) { + partial_sum[tid] += partial_sum[tid + s]; } __syncthreads(); } - // write result for this block to global mem - if (threadIdx.x == 0) { - result_data[threadIdx.x] = partial_sum[0]; + // Write block sum to global memory + if (tid == 0) { + result_data[blockIdx.x] = partial_sum[0]; } } -__host__ void sum_tensor_cuda(Tensor* tensor, float* result_data) { +__global__ void aux_final_sum_kernel(float* result_data, int size) { + __shared__ float partial_sum[THREADS_PER_BLOCK]; + int tid = threadIdx.x; + int i = blockIdx.x * blockDim.x + threadIdx.x; + + partial_sum[tid] = (i < size) ? result_data[i] : 0; + + __syncthreads(); + + // Perform final reduction + for (int s = blockDim.x / 2; s > 0; s >>= 1) { + if (tid < s) { + partial_sum[tid] += partial_sum[tid + s]; + } + __syncthreads(); + } + + // Write final result to global memory + if (tid == 0 && blockIdx.x == 0) { + result_data[0] = partial_sum[0]; + } +} + + + + +__host__ void sum_tensor_cuda(Tensor* tensor, float* result_data) { cudaMemcpy(result_data, tensor->data, tensor->size * sizeof(float), cudaMemcpyHostToDevice); - int number_of_blocks = (int)ceil(tensor->size / THREADS_PER_BLOCK); - sum_tensor_cuda_kernel<<>>(tensor->data, result_data); - sum_tensor_cuda_kernel<<<1, THREADS_PER_BLOCK>>>(result_data, result_data); + int num_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + + // First-level reduction + sum_tensor_cuda_kernel<<>>(tensor->data, result_data, tensor->size); + + // If necessary, perform multiple levels of reduction + while (num_blocks > 1) { + int num_blocks_next = (num_blocks + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + aux_final_sum_kernel<<>>(result_data, num_blocks); + num_blocks = num_blocks_next; + } cudaError_t error = cudaGetLastError(); if (error != cudaSuccess) { From 20fb09e6ea0902b6ca7771c1886bd825b5346b60 Mon Sep 17 00:00:00 2001 From: lucasdelimanogueira Date: Thu, 2 May 2024 14:47:36 -0300 Subject: [PATCH 2/2] Fix sum tensor reduce cuda v1 naive --- norch/csrc/cuda.cu | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index fb0ade6..9304207 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -4,6 +4,7 @@ #include #define THREADS_PER_BLOCK 128 +#define THREADS_PER_BLOCK_SUM 1024 #define TILE_SIZE 32 __host__ void cpu_to_cuda(Tensor* tensor) { @@ -60,7 +61,7 @@ __host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_da __global__ void sum_tensor_cuda_kernel(float* data, float* result_data, int size) { - __shared__ float partial_sum[THREADS_PER_BLOCK]; + __shared__ float partial_sum[THREADS_PER_BLOCK_SUM]; int tid = threadIdx.x; int i = blockIdx.x * blockDim.x + threadIdx.x; @@ -85,7 +86,7 @@ __global__ void sum_tensor_cuda_kernel(float* data, float* result_data, int size __global__ void aux_final_sum_kernel(float* result_data, int size) { - __shared__ float partial_sum[THREADS_PER_BLOCK]; + __shared__ float partial_sum[THREADS_PER_BLOCK_SUM]; int tid = threadIdx.x; int i = blockIdx.x * blockDim.x + threadIdx.x; @@ -114,15 +115,15 @@ __global__ void aux_final_sum_kernel(float* result_data, int size) { __host__ void sum_tensor_cuda(Tensor* tensor, float* result_data) { cudaMemcpy(result_data, tensor->data, tensor->size * sizeof(float), cudaMemcpyHostToDevice); - int num_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + int num_blocks = (tensor->size + THREADS_PER_BLOCK_SUM - 1) / THREADS_PER_BLOCK_SUM; // First-level reduction - sum_tensor_cuda_kernel<<>>(tensor->data, result_data, tensor->size); + sum_tensor_cuda_kernel<<>>(tensor->data, result_data, tensor->size); // If necessary, perform multiple levels of reduction while (num_blocks > 1) { - int num_blocks_next = (num_blocks + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; - aux_final_sum_kernel<<>>(result_data, num_blocks); + int num_blocks_next = (num_blocks + THREADS_PER_BLOCK_SUM - 1) / THREADS_PER_BLOCK_SUM; + aux_final_sum_kernel<<>>(result_data, num_blocks); num_blocks = num_blocks_next; }