Skip to content

CUDA: fix build error from ambiguous __half conversions in conv2d #15690

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

Merged
merged 5 commits into from
Sep 1, 2025
+3 −2
Merged

CUDA: fix build error from ambiguous __half conversions in conv2d #15690

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
2274449
CUDA: fix build error from ambiguous __half conversions in conv2d
qnixsynapse Aug 31, 2025
1d4e7ee
CUDA: Replace custom to_float helper with unified ggml_cuda_cast and …
qnixsynapse Aug 31, 2025
80a679b
CUDA: Add missing convert.cuh header
qnixsynapse Aug 31, 2025
a704126
CUDA: remove unnecessary extension in ggml_cuda_cast
qnixsynapse Aug 31, 2025
5251a40
CUDA: Address review comment, remove second type template argument
qnixsynapse Aug 31, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions ggml/src/ggml-cuda/conv2d.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,8 +94,8 @@ static __global__ void conv2d_kernel(const float * __restrict__ input,
const int64_t in_x = calculate_input_coord(out_x, kx, P.ST_X, P.DL_X, P.PD_X);

const float input_val = input[Layout::input_index(n, c_in, in_y, in_x, P)];
const float kernel_val = kernel[Layout::kernel_index(c_out, c_in, ky, kx, P)];
acc += (input_val * kernel_val);
const T kernel_val = kernel[Layout::kernel_index(c_out, c_in, ky, kx, P)];
acc += (input_val * ggml_cuda_cast<float, T>(kernel_val));
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
acc += (input_val * ggml_cuda_cast<float, T>(kernel_val));
acc += (input_val * ggml_cuda_cast<float>(kernel_val));

The second type can be inferred from the argument, I think this is more legible.

}
}
}
Expand Down
2 changes: 2 additions & 0 deletions ggml/src/ggml-cuda/convert.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,8 @@ template<typename dst_t, typename src_t>
__host__ __device__ inline dst_t ggml_cuda_cast(src_t x) {
if constexpr (std::is_same_v<dst_t, src_t>) {
return x;
} else if constexpr (std::is_same_v<dst_t, float> && std::is_same_v<src_t, half>) {
return __half2float(x);
} else if constexpr(std::is_same_v<dst_t, nv_bfloat16>) {
return __float2bfloat16(float(x));
} else if constexpr(std::is_same_v<src_t, nv_bfloat16>) {
Expand Down
Loading