Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

feat: cuda implementation for ggml_conv_transpose_1d #854

Open
wants to merge 11 commits into
base: master
Choose a base branch
from
Open

feat: cuda implementation for ggml_conv_transpose_1d #854

Changes from 1 commit
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
70de8b7
conv transpose 1d passing test for 1d input and kernel
balisujohn Jun 10, 2024
f6883de
working for different input and output channel counts, added test for…
balisujohn Jun 11, 2024
f35d3ec
initial draft appears to work with stride other than 1
balisujohn Jun 12, 2024
53a4fcf
working with all old and new conv1d tests
balisujohn Jun 13, 2024
f3bb758
added a test for large tensors
balisujohn Jun 13, 2024
7eff0ab
removed use cuda hardcoding
balisujohn Jun 13, 2024
152e04e
restored test-conv-transpose.c
balisujohn Jun 13, 2024
5d39cd4
removed unused arugments, and fixed bug where test failure would caus…
balisujohn Jun 14, 2024
2e7445e
fixed accumulator bug
balisujohn Jun 15, 2024
ed3b788
added test to test-backend-ops
balisujohn Jun 15, 2024
da3d0d1
fixed mistake
balisujohn Jun 15, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
Prev Previous commit
Next Next commit
fixed accumulator bug
  • Loading branch information
@balisujohn
balisujohn committed Jun 15, 2024
commit 2e7445e3a8437784b409dcd00375cc3cd2505573
6 changes: 3 additions & 3 deletions src/ggml-cuda/conv-transpose-1d.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ static __global__ void conv_transpose_1d_kernel(

int out_index = global_index / dst_ne0;

int accumulator = 0;
float accumulator = 0;

for (int c = 0; c < src0_ne2; c++)
{
Expand All @@ -34,8 +34,8 @@ static __global__ void conv_transpose_1d_kernel(
int weight_idx = idx - i*s0;


int kernel_weight = src0[kernel_offset + weight_idx];
int input_value = src1[input_offset+i];
float kernel_weight = src0[kernel_offset + weight_idx];
float input_value = src1[input_offset+i];

accumulator += kernel_weight * input_value;
}
Expand Down