aparler/PyNorch
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Fix cuda sum reduce

Browse source
This commit is contained in:
lucasdelimanogueira 2024-05-01 03:09:05 -03:00
parent a6da5bc9af
commit 9981bb8b6f
8 changed files with 12 additions and 10 deletions
Download patch file Download diff file Expand all files Collapse all files

BIN
build/cuda.cu.o
View file

Binary file not shown.

BIN
build/libtensor.so
View file

Binary file not shown.

BIN
build/tensor.o
View file

Binary file not shown.

BIN
norch/__pycache__/tensor.cpython-38.pyc
View file

Binary file not shown.

4
norch/csrc/cuda.cu
View file

@ -104,8 +104,8 @@ __host__ void sum_tensor_cuda(Tensor* tensor, float* result_data, int number_of_
int levelSize = number_of_blocks;
while (remaining > 1) {
int threads = (remaining + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
sum_block<<<1, threads, threads * sizeof(float)>>>(result_data, remaining);
remaining = (remaining + threadsPerBlock - 1) / threadsPerBlock;
aux_sum_block_kernel<<<1, threads, threads * sizeof(float)>>>(result_data, remaining);
remaining = (remaining + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
levelSize = remaining;
}

3
norch/csrc/cuda.h
View file

@ -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);
__host__ void sum_tensor_cuda(Tensor* tensor1, 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);
@ -36,4 +36,5 @@
#endif /* CUDA_KERNEL_H_ */

6
norch/csrc/tensor.cpp
View file

@ -171,9 +171,9 @@ extern "C" {
if (strcmp(tensor->device, "cuda") == 0) {
float* result_data;
int num_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
cudaMalloc((void **)&num_of_blocks, 1 * sizeof(float));
sum_tensor_cuda(tensor, result_data, int num_of_blocks);
int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
cudaMalloc((void**)&result_data, number_of_blocks * sizeof(float));
sum_tensor_cuda(tensor, result_data, number_of_blocks);
return create_tensor(result_data, shape, ndim, device);
}
else {

9
test.py
View file

@ -30,11 +30,12 @@ if __name__ == "__main__":
a = norch.Tensor([[1, 2], [1, 2], [1, 2]], requires_grad=True)#.to("cuda")
b = norch.Tensor([[1, 400, 3], [1, 2, 3]], requires_grad=True)
c = (b @ a) * 5
d = c.sum()
d.backward()
print(a.sum())
#c = (b @ a) * 5
#d = c.sum()
#d.backward()
print(a.grad)
#print(a.grad)
"""#print(a)
N = 1000