Skip to content

Windows: explicit instantiations repeat __restrict__, and GLM's NCCL defaults go through sys.c #1213

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

Closed
wesley4546 wants to merge 2 commits into from
Closed

Windows: explicit instantiations repeat __restrict__, and GLM's NCCL defaults go through sys.c #1213

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
0ca5628
cuda: explicit instantiations repeat their kernels' __restrict__, so …
wesley4546 Oct 23, 2026
786c2d6
glm cuda: openRank's NCCL defaults go through sys.c, so tensorfold-na…
wesley4546 Oct 23, 2026
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
15 changes: 10 additions & 5 deletions zig/kernels/cuda/exl3_block/dense3.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -379,24 +379,29 @@ template __global__ void dsv41_dense3::lanes_linear_kernel<4, 2, 4, 3>(const dsv
template __global__ void dsv41_dense3::lanes_linear_kernel<4, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<4, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<4, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_unpack_kernel<4, 2>(const uint32_t*, half*, int, int64_t, int64_t);
template __global__ void dsv41_dense3::lanes_unpack_kernel<4, 2>(const uint32_t* __restrict__, half* __restrict__, int,
int64_t, int64_t);
template __global__ void dsv41_dense3::lanes_linear_kernel<6, 2, 4, 3>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<6, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<6, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<6, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_unpack_kernel<6, 2>(const uint32_t*, half*, int, int64_t, int64_t);
template __global__ void dsv41_dense3::lanes_unpack_kernel<6, 2>(const uint32_t* __restrict__, half* __restrict__, int,
int64_t, int64_t);
template __global__ void dsv41_dense3::lanes_linear_kernel<8, 2, 4, 3>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<8, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<8, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<8, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_unpack_kernel<8, 2>(const uint32_t*, half*, int, int64_t, int64_t);
template __global__ void dsv41_dense3::lanes_unpack_kernel<8, 2>(const uint32_t* __restrict__, half* __restrict__, int,
int64_t, int64_t);
template __global__ void dsv41_dense3::lanes_linear_kernel<10, 2, 4, 3>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<10, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<10, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<10, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_unpack_kernel<10, 2>(const uint32_t*, half*, int, int64_t, int64_t);
template __global__ void dsv41_dense3::lanes_unpack_kernel<10, 2>(const uint32_t* __restrict__, half* __restrict__, int,
int64_t, int64_t);
template __global__ void dsv41_dense3::lanes_linear_kernel<12, 2, 4, 3>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<12, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<12, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_linear_kernel<12, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
template __global__ void dsv41_dense3::lanes_unpack_kernel<12, 2>(const uint32_t*, half*, int, int64_t, int64_t);
template __global__ void dsv41_dense3::lanes_unpack_kernel<12, 2>(const uint32_t* __restrict__, half* __restrict__, int,
int64_t, int64_t);
54 changes: 29 additions & 25 deletions zig/kernels/cuda/exl3_block/exl3_experts.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -245,8 +245,9 @@ __global__ void down_combine_kernel(const float* __restrict__ Z, const int* __re

// upstream grouped_kernel at mul1 for every (nt, warps, pf) x TF_RANGES, and dequant_kernel<2, K2> (tests)
#define TF_EXL3X_INST(NT, W, PF, LO, HI) template __global__ void tf_exl3x::grouped_kernel<2, NT, W, PF, LO, HI>( \
const half*, const half*, const int64_t*, const int64_t*, const int*, const int*, const int*, const int*, \
const int*, float*, int, int, int, int, int, int);
const half* __restrict__, const half* __restrict__, const int64_t* __restrict__, const int64_t* __restrict__, \
const int* __restrict__, const int* __restrict__, const int* __restrict__, const int* __restrict__, \
const int* __restrict__, float* __restrict__, int, int, int, int, int, int);
TF_EXL3X_INST(8, 4, 1, 8, 8)
TF_EXL3X_INST(8, 4, 1, 2, 10)
TF_EXL3X_INST(8, 4, 1, 2, 16)
Expand All @@ -256,28 +257,31 @@ TF_EXL3X_INST(8, 4, 2, 2, 16)
TF_EXL3X_INST(4, 4, 2, 8, 8)
TF_EXL3X_INST(4, 4, 2, 2, 10)
TF_EXL3X_INST(4, 4, 2, 2, 16)
template __global__ void tf_exl3x::dequant_kernel<2, 2>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 3>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 4>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 5>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 6>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 7>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 8>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 9>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 10>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 11>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 12>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 13>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 14>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 15>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 16>(const uint32_t*, half*, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 2>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 3>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 4>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 5>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 6>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 7>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 8>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 9>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 10>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 11>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 12>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 13>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 14>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 15>(const uint32_t* __restrict__, half* __restrict__, int, int);
template __global__ void tf_exl3x::dequant_kernel<2, 16>(const uint32_t* __restrict__, half* __restrict__, int, int);
// down_combine_kernel at both outputs the binding takes (fp32, bf16: TF_DSV41_MOE_BF16)
template __global__ void tf_exl3_experts::down_combine_kernel<float>(const float*, const int*, const half*, float*, \
const float*, float*, int, int, int, int, int);
template __global__ void tf_exl3_experts::down_combine_kernel<__nv_bfloat16>(const float*, const int*, const half*, float*, \
const float*, __nv_bfloat16*, int, int, int, int, int);
template __global__ void tf_exl3_experts::down_combine_kernel<float>(const float* __restrict__, \
const int* __restrict__, const half* __restrict__, float* __restrict__, const float* __restrict__, \
float* __restrict__, int, int, int, int, int);
template __global__ void tf_exl3_experts::down_combine_kernel<__nv_bfloat16>(const float* __restrict__, \
const int* __restrict__, const half* __restrict__, float* __restrict__, const float* __restrict__, \
__nv_bfloat16* __restrict__, int, int, int, int, int);
// rot_in_kernel for both inputs the binding takes
template __global__ void tf_exl3_experts::rot_in_kernel<__nv_bfloat16>(const __nv_bfloat16*, int, const int*, const half*, const half*, half*, \
half*, int, int, int);
template __global__ void tf_exl3_experts::rot_in_kernel<half>(const half*, int, const int*, const half*, const half*, half*, \
half*, int, int, int);
template __global__ void tf_exl3_experts::rot_in_kernel<__nv_bfloat16>(const __nv_bfloat16* __restrict__, int, \
const int* __restrict__, const half* __restrict__, const half* __restrict__, half* __restrict__, \
half* __restrict__, int, int, int);
template __global__ void tf_exl3_experts::rot_in_kernel<half>(const half* __restrict__, int, const int* __restrict__, \
const half* __restrict__, const half* __restrict__, half* __restrict__, half* __restrict__, int, int, int);
10 changes: 6 additions & 4 deletions zig/kernels/cuda/exl3_block/lin2.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -467,10 +467,12 @@ __global__ void __launch_bounds__(WK * 32, MINB) lanesc_kernel(const dsv41_x3seg
template __global__ void dsv41_lin2::lanes2_kernel<K2, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
LIN2_LANES(4) LIN2_LANES(6) LIN2_LANES(8) LIN2_LANES(10) LIN2_LANES(12)
#define LIN2_LINEAR(K2, WK) \
template __global__ void dsv41_lin2::linear2_kernel<K2, 2, WK, true>(const half*, const uint32_t*, long long, \
long long, const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*); \
template __global__ void dsv41_lin2::linear2_kernel<K2, 2, WK, false>(const half*, const uint32_t*, long long, \
long long, const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
template __global__ void dsv41_lin2::linear2_kernel<K2, 2, WK, true>(const half* __restrict__, \
const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__); \
template __global__ void dsv41_lin2::linear2_kernel<K2, 2, WK, false>(const half* __restrict__, \
const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
#define LIN2_LINEAR_W(K2) LIN2_LINEAR(K2, 2) LIN2_LINEAR(K2, 4) LIN2_LINEAR(K2, 8)
LIN2_LINEAR_W(2) LIN2_LINEAR_W(3) LIN2_LINEAR_W(4) LIN2_LINEAR_W(5) LIN2_LINEAR_W(6) LIN2_LINEAR_W(7)
LIN2_LINEAR_W(8) LIN2_LINEAR_W(10) LIN2_LINEAR_W(12) LIN2_LINEAR_W(14) LIN2_LINEAR_W(16)
Expand Down
Loading