diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index 91b349c..ffb93af 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -57,7 +57,7 @@ __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) { +_global__ void sum_tensor_cuda_kernel(float* data, float* result_data, int size) { extern __shared__ float sdata[]; unsigned int tid = threadIdx.x; @@ -73,7 +73,7 @@ __global__ void sum_tensor_cuda_kernel(float* data, float* result_data, int size // Reduce within block using shared memory for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { - if (tid < s) { + if (tid < s && i + s < size) { sdata[tid] += sdata[tid + s]; } __syncthreads(); diff --git a/norch/tensor.py b/norch/tensor.py index 0a32a61..74be89f 100644 --- a/norch/tensor.py +++ b/norch/tensor.py @@ -20,7 +20,6 @@ class Tensor: if data != None: data, shape = self.flatten(data) - # Adjust the path to the shared library self.data_ctype = (ctypes.c_float * len(data))(*data) self.shape_ctype = (ctypes.c_int * len(shape))(*shape) self.ndim_ctype = ctypes.c_int(len(shape))