From 8e11d7537a0a6b91a48688c2a3263112257e3a62 Mon Sep 17 00:00:00 2001 From: lucasdelimanogueira Date: Thu, 2 May 2024 14:58:35 -0300 Subject: [PATCH 1/3] Fix sum tensor reduce cuda v1 naive --- norch/csrc/cuda.cu | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index 9304207..d1b6a49 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -6,6 +6,7 @@ #define THREADS_PER_BLOCK 128 #define THREADS_PER_BLOCK_SUM 1024 #define TILE_SIZE 32 +#define SHMEM_SIZE THREADS_PER_BLOCK_SUM * sizeof(float) __host__ void cpu_to_cuda(Tensor* tensor) { @@ -61,7 +62,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_SUM]; + __shared__ float partial_sum[SHMEM_SIZE]; int tid = threadIdx.x; int i = blockIdx.x * blockDim.x + threadIdx.x; @@ -86,7 +87,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_SUM]; + __shared__ float partial_sum[SHMEM_SIZE]; int tid = threadIdx.x; int i = blockIdx.x * blockDim.x + threadIdx.x; @@ -137,6 +138,8 @@ __host__ void sum_tensor_cuda(Tensor* tensor, float* result_data) { } + + __global__ void sub_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size) { int i = blockIdx.x * blockDim.x + threadIdx.x; From 04ada494a515396c1e353ea528dd96492d5809e2 Mon Sep 17 00:00:00 2001 From: lucasdelimanogueira Date: Thu, 2 May 2024 15:15:56 -0300 Subject: [PATCH 2/3] Fix reduce sum cuda large arrays --- norch/csrc/cuda.cu | 30 +----------------------------- 1 file changed, 1 insertion(+), 29 deletions(-) diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index d1b6a49..45f5f28 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -85,34 +85,6 @@ __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[SHMEM_SIZE]; - - 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); @@ -124,7 +96,7 @@ __host__ void sum_tensor_cuda(Tensor* tensor, float* result_data) { // If necessary, perform multiple levels of reduction while (num_blocks > 1) { int num_blocks_next = (num_blocks + THREADS_PER_BLOCK_SUM - 1) / THREADS_PER_BLOCK_SUM; - aux_final_sum_kernel<<>>(result_data, num_blocks); + sum_tensor_cuda_kernel<<>>(result_data, result_data, num_blocks); num_blocks = num_blocks_next; } From 976f932007cf9d23e240dce0780894a8ce62046b Mon Sep 17 00:00:00 2001 From: lucasdelimanogueira Date: Thu, 2 May 2024 15:18:14 -0300 Subject: [PATCH 3/3] small fix threads per block sum --- norch/csrc/cuda.cu | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index 45f5f28..bf413e3 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -4,7 +4,6 @@ #include #define THREADS_PER_BLOCK 128 -#define THREADS_PER_BLOCK_SUM 1024 #define TILE_SIZE 32 #define SHMEM_SIZE THREADS_PER_BLOCK_SUM * sizeof(float) @@ -88,15 +87,15 @@ __global__ void sum_tensor_cuda_kernel(float* data, 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_SUM - 1) / THREADS_PER_BLOCK_SUM; + 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); + 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_SUM - 1) / THREADS_PER_BLOCK_SUM; - sum_tensor_cuda_kernel<<>>(result_data, result_data, num_blocks); + int num_blocks_next = (num_blocks + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + sum_tensor_cuda_kernel<<>>(result_data, result_data, num_blocks); num_blocks = num_blocks_next; }