diff --git a/build/cuda.cu.o b/build/cuda.cu.o index 24a944b..8bbf51d 100644 Binary files a/build/cuda.cu.o and b/build/cuda.cu.o differ diff --git a/build/libtensor.so b/build/libtensor.so index 55b04fb..8dc7cf6 100755 Binary files a/build/libtensor.so and b/build/libtensor.so differ diff --git a/build/tensor.o b/build/tensor.o index 915cd04..13c474d 100644 Binary files a/build/tensor.o and b/build/tensor.o differ diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index 1dd5ed2..c24e48a 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -3,6 +3,10 @@ #include #include +#define THREADS_PER_BLOCK 128 +#define TILE_SIZE 32 +#define SHMEM_SIZE THREADS_PER_BLOCK * sizeof(float) + __host__ void cpu_to_cuda(Tensor* tensor) { float* data_tmp; @@ -55,60 +59,40 @@ __host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_da cudaDeviceSynchronize(); } -__global__ void sum_tensor_cuda_kernel(float* data, float* result_data, int size) { - extern __shared__ float sdata[]; - - unsigned int tid = threadIdx.x; - unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; - - sdata[tid] = (i < size) ? data[i] : 0; - __syncthreads(); - - for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { - if (tid < s) { - sdata[tid] += sdata[tid + s]; - } - __syncthreads(); - } - - if (tid == 0) { - result_data[blockIdx.x] = sdata[0]; - } -} __global__ void sum_tensor_cuda_kernel(float* data, float* result_data) { - __shared__ int sdata[THREADS_PER_BLOCK]; + __shared__ int partial_sum[SHMEM_SIZE]; // 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; - sdata[threadIdx.x] = data[i]; + partial_sum[threadIdx.x] = data[i]; __syncthreads(); // do reduction in shared mem for (int s=1; s < blockDim.x; s *=2) { - int index = 2 * s * threadIdx.x;; - if (index < blockDim.x) - { - sdata[index] += sdata[index + s]; + if (threadIdx.x % (2 * s) == 0) { + partial_sum[threadIdx.x] += partial_sum[threadIdx.x + s]; } __syncthreads(); } // write result for this block to global mem - if (threadIdx.x == 0) - atomicAdd(result_data, sdata[0]); + if (threadIdx.x == 0) { + result_data[threadIdx.x] = partial_sum[0]; + } } __host__ void sum_tensor_cuda(Tensor* tensor, float* result_data) { - number_of_blocks = tensor->size / THREADS_PER_BLOCK; - sum_tensor_cuda_kernel<<>>(tensor->data, result_data, tensor->size); + 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); cudaError_t error = cudaGetLastError(); if (error != cudaSuccess) { diff --git a/norch/csrc/cuda.h b/norch/csrc/cuda.h index 7825534..48c87b8 100644 --- a/norch/csrc/cuda.h +++ b/norch/csrc/cuda.h @@ -7,8 +7,8 @@ __global__ void add_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size); __host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data); - __host__ void sum_tensor_cuda(Tensor* tensor, float* result_data, int number_of_blocks); - __global__ void sum_tensor_cuda_kernel(float* data, float* result_data, int size); + __global__ void sum_tensor_cuda_kernel(float* data, float* result_data); + __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); __host__ void sub_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data); diff --git a/norch/csrc/tensor.h b/norch/csrc/tensor.h index b09620d..aa8d3e9 100644 --- a/norch/csrc/tensor.h +++ b/norch/csrc/tensor.h @@ -1,9 +1,6 @@ #ifndef TENSOR_H #define TENSOR_H -#define THREADS_PER_BLOCK 128 -#define TILE_SIZE 32 - typedef struct { float* data; int* strides;