zhangyiss/mlx
1
0
Fork 0
mirror of https://github.com/ml-explore/mlx.git synced 2025-09-01 12:49:44 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Add dynamic shared memory and improve qmm

Browse Source
This commit is contained in:
Angelos Katharopoulos
2025-07-22 23:36:53 -07:00
parent 700f7dcf01
commit 903b40627c
28 changed files with 273 additions and 82 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
mlx/backend/cuda/arg_reduce.cu
View File

@@ -166,6 +166,7 @@ void ArgReduce::eval_gpu(const std::vector<array>& inputs, array& out) {
kernel,
num_blocks,
block_dim(),
0,
in.data<T>(),
out.data<uint32_t>(),
out.size(),

3
mlx/backend/cuda/binary.cu
View File

@@ -219,6 +219,7 @@ void binary_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
a.data<InType>(),
b.data<InType>(),
out.data<OutType>(),
@@ -235,6 +236,7 @@ void binary_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
a.data<InType>(),
b.data<InType>(),
out.data<OutType>(),
@@ -269,6 +271,7 @@ void binary_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
a.data<InType>(),
b.data<InType>(),
out.data<OutType>(),

3
mlx/backend/cuda/binary_two.cu
View File

@@ -239,6 +239,7 @@ void binary_two_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
a.data<InType>(),
b.data<InType>(),
out_a.data<OutType>(),
@@ -256,6 +257,7 @@ void binary_two_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
a.data<InType>(),
b.data<InType>(),
out_a.data<OutType>(),
@@ -291,6 +293,7 @@ void binary_two_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
a.data<InType>(),
b.data<InType>(),
out_a.data<OutType>(),

2
mlx/backend/cuda/compiled.cpp
View File

@@ -295,7 +295,7 @@ void Compiled::eval_gpu(
auto kernel = mod.get_kernel(kernel_name);
auto [num_blocks, block_dims] =
get_launch_args(kernel, outputs[0], large, work_per_thread);
encoder.add_kernel_node(kernel, num_blocks, block_dims, args.args());
encoder.add_kernel_node(kernel, num_blocks, block_dims, 0, args.args());
}
} // namespace mlx::core

1
mlx/backend/cuda/copy/copy_contiguous.cu
View File

@@ -82,6 +82,7 @@ void copy_contiguous(
kernel,
num_blocks,
block_dims,
0,
in.data<InType>() + in_offset,
out.data<OutType>() + out_offset,
out.data_size());

2
mlx/backend/cuda/copy/copy_general.cu
View File

@@ -79,6 +79,7 @@ void copy_general(
kernel,
num_blocks,
block_dims,
0,
in_ptr,
out_ptr,
data_size,
@@ -94,6 +95,7 @@ void copy_general(
kernel,
num_blocks,
block_dims,
0,
in_ptr,
out_ptr,
data_size,

2
mlx/backend/cuda/copy/copy_general_dynamic.cu
View File

@@ -82,6 +82,7 @@ void copy_general_dynamic(
kernel,
num_blocks,
block_dims,
0,
in_ptr,
out_ptr,
out.size(),
@@ -99,6 +100,7 @@ void copy_general_dynamic(
kernel,
num_blocks,
block_dims,
0,
in_ptr,
out_ptr,
out.size(),

2
mlx/backend/cuda/copy/copy_general_input.cu
View File

@@ -71,6 +71,7 @@ void copy_general_input(
kernel,
num_blocks,
block_dims,
0,
in_ptr,
out_ptr,
out.size(),
@@ -85,6 +86,7 @@ void copy_general_input(
kernel,
num_blocks,
block_dims,
0,
in_ptr,
out_ptr,
out.size(),

4
mlx/backend/cuda/device.cpp
View File

@@ -215,12 +215,14 @@ void CommandEncoder::add_kernel_node(
void* func,
dim3 grid_dim,
dim3 block_dim,
uint32_t smem_bytes,
void** params) {
cudaKernelNodeParams kernel_params = {0};
kernel_params.func = func;
kernel_params.gridDim = grid_dim;
kernel_params.blockDim = block_dim;
kernel_params.kernelParams = params;
kernel_params.sharedMemBytes = smem_bytes;
cudaGraphNode_t node;
CHECK_CUDA_ERROR(
cudaGraphAddKernelNode(&node, graph_, NULL, 0, &kernel_params));
@@ -231,6 +233,7 @@ void CommandEncoder::add_kernel_node(
CUfunction func,
dim3 grid_dim,
dim3 block_dim,
uint32_t smem_bytes,
void** params) {
CUDA_KERNEL_NODE_PARAMS kernel_params = {0};
kernel_params.func = func;
@@ -241,6 +244,7 @@ void CommandEncoder::add_kernel_node(
kernel_params.blockDimY = block_dim.y;
kernel_params.blockDimZ = block_dim.z;
kernel_params.kernelParams = params;
kernel_params.sharedMemBytes = smem_bytes;
CUgraphNode node;
CHECK_CUDA_ERROR(
cuGraphAddKernelNode(&node, graph_, NULL, 0, &kernel_params));

19
mlx/backend/cuda/device.h
View File

@@ -45,25 +45,34 @@ class CommandEncoder {
void set_output_array(const array& arr);
template <typename F, typename... Params>
void
add_kernel_node(F* func, dim3 grid_dim, dim3 block_dim, Params&&... params) {
void add_kernel_node(
F* func,
dim3 grid_dim,
dim3 block_dim,
uint32_t smem_bytes,
Params&&... params) {
constexpr size_t num = sizeof...(Params);
void* ptrs[num];
size_t i = 0;
([&](auto&& p) { ptrs[i++] = static_cast<void*>(&p); }(
std::forward<Params>(params)),
...);
add_kernel_node((void*)func, grid_dim, block_dim, ptrs);
add_kernel_node((void*)func, grid_dim, block_dim, smem_bytes, ptrs);
}
void add_kernel_node(
CUfunction func,
dim3 grid_dim,
dim3 block_dim,
uint32_t smem_bytes,
void** params);
void
add_kernel_node(void* func, dim3 grid_dim, dim3 block_dim, void** params);
void add_kernel_node(
void* func,
dim3 grid_dim,
dim3 block_dim,
uint32_t smem_bytes,
void** params);
void add_temporary(const array& arr) {
temporaries_.push_back(arr.data_shared_ptr());

8
mlx/backend/cuda/indexing.cpp
View File

@@ -129,7 +129,7 @@ void Gather::eval_gpu(const std::vector<array>& inputs, array& out) {
auto kernel = mod.get_kernel(kernel_name);
auto [num_blocks, block_dims] = get_launch_args(kernel, out, large);
encoder.add_kernel_node(kernel, num_blocks, block_dims, args.args());
encoder.add_kernel_node(kernel, num_blocks, block_dims, 0, args.args());
}
void Scatter::eval_gpu(const std::vector<array>& inputs, array& out) {
@@ -230,7 +230,7 @@ void Scatter::eval_gpu(const std::vector<array>& inputs, array& out) {
encoder.set_output_array(out);
auto kernel = mod.get_kernel(kernel_name);
auto [num_blocks, block_dims] = get_launch_args(kernel, upd, large);
encoder.add_kernel_node(kernel, num_blocks, block_dims, args.args());
encoder.add_kernel_node(kernel, num_blocks, block_dims, 0, args.args());
}
void GatherAxis::eval_gpu(const std::vector<array>& inputs, array& out) {
@@ -318,7 +318,7 @@ void GatherAxis::eval_gpu(const std::vector<array>& inputs, array& out) {
encoder.set_output_array(out);
auto kernel = mod.get_kernel(kernel_name);
auto [num_blocks, block_dims] = get_launch_args(kernel, idx, large);
encoder.add_kernel_node(kernel, num_blocks, block_dims, args.args());
encoder.add_kernel_node(kernel, num_blocks, block_dims, 0, args.args());
}
void ScatterAxis::eval_gpu(const std::vector<array>& inputs, array& out) {
@@ -422,7 +422,7 @@ void ScatterAxis::eval_gpu(const std::vector<array>& inputs, array& out) {
encoder.set_output_array(out);
auto kernel = mod.get_kernel(kernel_name);
auto [num_blocks, block_dims] = get_launch_args(kernel, idx, large);
encoder.add_kernel_node(kernel, num_blocks, block_dims, args.args());
encoder.add_kernel_node(kernel, num_blocks, block_dims, 0, args.args());
}
} // namespace mlx::core

2
mlx/backend/cuda/layer_norm.cu
View File

@@ -266,6 +266,7 @@ void LayerNorm::eval_gpu(
kernel,
n_rows,
block_dim(),
0,
x.data<DataType>(),
w.data<DataType>(),
b.data<DataType>(),
@@ -378,6 +379,7 @@ void LayerNormVJP::eval_gpu(
kernel,
n_rows,
block_dim(),
0,
x.data<DataType>(),
w.data<DataType>(),
g.data<DataType>(),

1
mlx/backend/cuda/logsumexp.cu
View File

@@ -151,6 +151,7 @@ void LogSumExp::eval_gpu(const std::vector<array>& inputs, array& out) {
kernel,
n_rows,
block_dim(),
0,
in.data<DataType>(),
out.data<DataType>(),
axis_size);

34
mlx/backend/cuda/matmul/mma.cuh
View File

@@ -6,8 +6,9 @@
namespace mlx::core::cu {
template <typename TileAccum, typename Tile>
__device__ inline void mma(TileAccum& C, Tile& A, Tile& B) {}
template <typename U, typename T>
__device__ inline void
mma_t(Tile16x16<U>& C, Tile16x16<T>& A, Tile16x16<T>& B) {}
/**
* Multiply the 16x16 bfloat16 tiles and accumulate the result in one 16x16
@@ -15,7 +16,7 @@ __device__ inline void mma(TileAccum& C, Tile& A, Tile& B) {}
*
* We actually perform C += A @ B.T
*/
__device__ inline void mma(
__device__ inline void mma_t(
Tile16x16<float>& C,
Tile16x16<__nv_bfloat16>& A,
Tile16x16<__nv_bfloat16>& B) {
@@ -77,4 +78,31 @@ __device__ inline void mma(
"f"(C.values[3].y));
}
/**
* Multiply larger register tiles by delegating to mma_t.
*/
template <typename U, typename T, int M, int N, int K>
__device__ inline void mma_t(
RegisterTile<U, M, N>& C,
RegisterTile<T, M, K>& A,
RegisterTile<T, N, K>& B) {
constexpr int TILES_M = RegisterTile<T, M, K>::TILES_Y;
constexpr int TILES_K = RegisterTile<T, M, K>::TILES_X;
constexpr int TILES_N = RegisterTile<T, N, K>::TILES_Y;
MLX_UNROLL
for (int k = 0; k < TILES_K; k++) {
MLX_UNROLL
for (int m = 0; m < TILES_M; m++) {
MLX_UNROLL
for (int n = 0; n < TILES_N; n++) {
mma_t(
C.data[m * TILES_N + n],
A.data[m * TILES_K + k],
B.data[n * TILES_K + k]);
}
}
}
}
} // namespace mlx::core::cu

132
mlx/backend/cuda/matmul/tiles.cuh
View File

@@ -2,6 +2,8 @@
#pragma once
#define MLX_UNROLL _Pragma("unroll")
namespace mlx::core::cu {
// Map types to their vector of 2 type float -> float2, double -> double2 etc
@@ -41,9 +43,10 @@ struct Tile16x16 {
T2 values[4];
__device__ inline void clear() {
__device__ inline void fill(T v) {
T2 v2 = {v, v};
for (int i = 0; i < 4; i++) {
values[i] = static_cast<T2>(0);
values[i] = v2;
}
}
@@ -121,6 +124,57 @@ struct Tile16x16 {
}
};
/**
* A simple container of multiple Tile16x16.
*
* Provides utility functions for loading and manipulating collections of basic
* tiles.
*/
template <typename T, int ROWS_, int COLS_>
struct RegisterTile {
static constexpr int ROWS = ROWS_;
static constexpr int COLS = COLS_;
static constexpr int TILES_X = COLS / 16;
static constexpr int TILES_Y = ROWS / 16;
Tile16x16<T> data[TILES_X * TILES_Y];
__device__ inline void fill(T v) {
MLX_UNROLL
for (int i = 0; i < TILES_Y; i++) {
MLX_UNROLL
for (int j = 0; j < TILES_X; j++) {
data[i * TILES_X + j].fill(v);
}
}
}
template <typename Tile>
__device__ inline void
load(Tile& tile, uint32_t base_address, int row, int col) {
MLX_UNROLL
for (int i = 0; i < TILES_Y; i++) {
MLX_UNROLL
for (int j = 0; j < TILES_X; j++) {
data[i * TILES_X + j].load(
tile.loc(base_address, row + i * 16, col + j * 16));
}
}
}
template <typename U>
__device__ inline void store_global(U* x, int N, int row, int col) {
MLX_UNROLL
for (int i = 0; i < TILES_Y; i++) {
MLX_UNROLL
for (int j = 0; j < TILES_X; j++) {
data[i * TILES_X + j].store_global(
x + (row + i * 16) * N + col + j * 16, N);
}
}
}
};
template <typename T, int ROWS_, int COLS_>
struct SharedTile {
static constexpr int ROWS = ROWS_;
@@ -198,7 +252,7 @@ struct SharedTile {
} else if constexpr (sizeof(T) * N == 16) {
store(*(reinterpret_cast<float4*>(&v[0])), row, col);
} else {
#pragma unroll
MLX_UNROLL
for (int i = 0; i < N; i++) {
*ptr(data, row, col + i) = v[i];
}
@@ -206,6 +260,10 @@ struct SharedTile {
}
};
/**
* Load the tile from global memory by loading 16 bytes at a time and storing
* them immediately.
*/
template <int NUM_WARPS, typename T, typename Tile>
__device__ inline void load(Tile& tile, const T* x, int N) {
constexpr int NUM_THREADS = NUM_WARPS * 32;
@@ -220,7 +278,7 @@ __device__ inline void load(Tile& tile, const T* x, int N) {
x += row * N + col * ELEMENTS_PER_LOAD;
#pragma unroll
MLX_UNROLL
for (int i = 0; i < NUM_LOADS_PER_THREAD; i++) {
float4 tmp;
tmp = *(reinterpret_cast<const float4*>(&x[i * STEP_ROWS * N]));
@@ -228,4 +286,70 @@ __device__ inline void load(Tile& tile, const T* x, int N) {
}
}
/**
* Copy 16 bytes from the globale memory address pointed to by x to the smem
* address pointed to by row_address.
*
* A simple wrapper over the PTX.
*/
template <typename T>
__device__ inline void cp_async_16(uint32_t row_address, const T* x) {
asm volatile(
"cp.async.ca.shared::cta.global [%0], [%1], 16;\n" ::"r"(row_address),
"l"(reinterpret_cast<const int4*>(x)));
}
/**
* Submit all the previous async copies to be executed.
*/
__device__ inline void cp_async_commit() {
asm volatile("cp.async.commit_group;\n" ::);
}
/**
* Wait for all the async copies to finish.
*/
__device__ inline void cp_async_wait_all() {
asm volatile("cp.async.wait_all;\n" ::);
}
/**
* The asynchronous equivalent of load.
*
* Loads the tile from global memory by submitting a bunch of async copy
* instructions. The copy won't start until commit is called and we don't have
* a guarantee it will finish until wait is called.
*
* It should be used as follows
*
* load(...)
* load(...)
* cp_async_commit()
* do_other_stuff()
* cp_async_wait_all()
* do_stuff_with_shmem()
*/
template <int NUM_WARPS, typename T, typename Tile>
__device__ inline void
load_async(Tile& tile, uint32_t base_address, const T* x, int N) {
constexpr int NUM_THREADS = NUM_WARPS * 32;
constexpr int ELEMENTS_PER_LOAD = sizeof(float4) / sizeof(T);
constexpr int NUM_LOADS = Tile::NUMEL / ELEMENTS_PER_LOAD;
constexpr int NUM_LOADS_PER_THREAD = NUM_LOADS / NUM_THREADS;
constexpr int NUM_LOADS_PER_ROW = Tile::COLS / ELEMENTS_PER_LOAD;
constexpr int STEP_ROWS = NUM_THREADS / NUM_LOADS_PER_ROW;
const int row = threadIdx.x / NUM_LOADS_PER_ROW;
const int col = threadIdx.x % NUM_LOADS_PER_ROW;
x += row * N + col * ELEMENTS_PER_LOAD;
MLX_UNROLL
for (int i = 0; i < NUM_LOADS_PER_THREAD; i++) {
cp_async_16(
tile.loc(base_address, row + i * STEP_ROWS, col * ELEMENTS_PER_LOAD),
x + i * STEP_ROWS * N);
}
}
} // namespace mlx::core::cu

2
mlx/backend/cuda/quantized/affine_quantize.cu
View File

@@ -261,6 +261,7 @@ void affine_quantize(
kernel,
num_blocks,
block_dims,
0,
w.data<T>(),
wq.data<uint8_t>(),
scales.data<T>(),
@@ -316,6 +317,7 @@ void affine_dequantize(
kernel,
num_blocks,
block_dims,
0,
wq.data<uint8_t>(),
scales.data<T>(),
biases.data<T>(),

110
mlx/backend/cuda/quantized/qmm.cu
View File

@@ -36,13 +36,13 @@ __device__ inline void load_quantized(
scales += row * Ng + col * ELEMENTS_PER_LOAD / group_size;
biases += row * Ng + col * ELEMENTS_PER_LOAD / group_size;
#pragma unroll
MLX_UNROLL
for (int i = 0; i < NUM_LOADS_PER_THREAD; i++) {
T vs[ELEMENTS_PER_LOAD];
uint32_t w = *reinterpret_cast<const uint32_t*>(x + i * STEP_ROWS * Nx);
T s = scales[i * STEP_ROWS * Ng];
T b = biases[i * STEP_ROWS * Ng];
#pragma unroll
MLX_UNROLL
for (int j = 0; j < ELEMENTS_PER_LOAD; j++) {
vs[j] = static_cast<T>((w >> (j * bits)) & MASK) * s + b;
}
@@ -51,7 +51,7 @@ __device__ inline void load_quantized(
}
template <typename T, int BM, int BN, int BK, int group_size, int bits>
__global__ void qmm(
__global__ void qmm_t_aligned(
const T* x,
const uint8_t* w,
const T* scales,
@@ -60,24 +60,27 @@ __global__ void qmm(
int M,
int N,
int K) {
constexpr int NUM_WARPS = 4;
constexpr int WARP_M = (BM / 16) / (NUM_WARPS / 2);
constexpr int WARP_N = (BN / 16) / (NUM_WARPS / 2);
constexpr int WARP_K = BK / 16;
constexpr int WARP_STEP_M = WARP_M * 16;
constexpr int WARP_STEP_N = WARP_N * 16;
constexpr int WARPS_M = 2;
constexpr int WARPS_N = 4;
constexpr int NUM_WARPS = WARPS_M * WARPS_N;
constexpr int WARP_STEP_M = BM / WARPS_M;
constexpr int WARP_STEP_N = BN / WARPS_N;
const int warpid = threadIdx.x / 32;
const int laneid = threadIdx.x % 32;
const int offset_m = (warpid / 2) * WARP_STEP_M;
const int offset_n = (warpid % 2) * WARP_STEP_N;
const int wm = warpid / WARPS_N;
const int wn = warpid % WARPS_N;
const int offset_m = wm * WARP_STEP_M;
const int offset_n = wn * WARP_STEP_N;
__shared__ SharedTile<T, BM, BK> xs;
__shared__ SharedTile<T, BN, BK> ws;
extern __shared__ char shmem[];
SharedTile<T, BM, BK>(&xs)[1] = *(SharedTile<T, BM, BK>(*)[1])(&shmem[0]);
SharedTile<T, BN, BK>(&ws)[1] =
*(SharedTile<T, BN, BK>(*)[1])(&shmem[1 * sizeof(T) * BM * BK]);
Tile16x16<float> C[WARP_M * WARP_N];
Tile16x16<T> A[WARP_M];
Tile16x16<T> B[WARP_N];
RegisterTile<float, BM / WARPS_M, BN / WARPS_N> C;
RegisterTile<T, BM / WARPS_M, 16> A;
RegisterTile<T, BN / WARPS_N, 16> B;
x += blockIdx.y * BM * K;
w += blockIdx.x * BN * K / get_pack_factor<bits>();
@@ -85,59 +88,43 @@ __global__ void qmm(
biases += blockIdx.x * BN * K / group_size;
y += blockIdx.y * BM * N + blockIdx.x * BN;
#pragma unroll
for (int i = 0; i < WARP_M * WARP_N; i++) {
C[i].clear();
}
C.fill(0);
uint32_t base_addr_xs = __cvta_generic_to_shared(&xs.data[0]);
uint32_t base_addr_ws = __cvta_generic_to_shared(&ws.data[0]);
int tic = 0;
uint32_t base_addr_xs[1], base_addr_ws[1];
base_addr_xs[0] = __cvta_generic_to_shared(&xs[0].data[0]);
base_addr_ws[0] = __cvta_generic_to_shared(&ws[0].data[0]);
for (int k_block = 0; k_block < K; k_block += BK) {
load<NUM_WARPS>(xs, x + k_block, K);
load_async<NUM_WARPS>(xs[tic], base_addr_xs[tic], x + k_block, K);
cp_async_commit();
// load<NUM_WARPS>(xs[tic], x + k_block, K);
load_quantized<NUM_WARPS, group_size, bits>(
ws,
ws[tic],
w + k_block / get_pack_factor<bits>(),
scales + k_block / group_size,
biases + k_block / group_size,
K);
cp_async_wait_all();
__syncthreads();
#pragma unroll
for (int k = 0; k < WARP_K; k++) {
#pragma unroll
for (int i = 0; i < WARP_M; i++) {
A[i].load(xs.loc(
base_addr_xs,
offset_m + i * 16 + laneid % 16,
k * 16 + laneid / 16 * 8));
}
#pragma unroll
for (int i = 0; i < WARP_N; i++) {
B[i].load(ws.loc(
base_addr_ws,
offset_n + i * 16 + laneid % 16,
k * 16 + laneid / 16 * 8));
}
#pragma unroll
for (int i = 0; i < WARP_M; i++) {
#pragma unroll
for (int j = 0; j < WARP_N; j++) {
mma(C[i * WARP_N + j], A[i], B[j]);
}
}
MLX_UNROLL
for (int k = 0; k < BK / 16; k++) {
A.load(
xs[tic],
base_addr_xs[tic],
offset_m + laneid % 16,
k * 16 + laneid / 16 * 8);
B.load(
ws[tic],
base_addr_ws[tic],
offset_n + laneid % 16,
k * 16 + laneid / 16 * 8);
mma_t(C, A, B);
}
}
#pragma unroll
for (int i = 0; i < WARP_M; i++) {
#pragma unroll
for (int j = 0; j < WARP_N; j++) {
C[i * WARP_N + j].store_global(
y + (offset_m + i * 16) * N + offset_n + j * 16, N);
}
}
C.store_global(y, N, offset_m, offset_n);
}
} // namespace cu
@@ -160,18 +147,19 @@ void qmm(
dispatch_groups(group_size_, [&](auto group_size) {
dispatch_bits(bits_, [&](auto bits) {
using DataType = cuda_type_t<MLX_GET_TYPE(type_tag)>;
constexpr int BM = 64;
constexpr int BN = 64;
constexpr int BM = 128;
constexpr int BN = 128;
constexpr int BK = 32;
auto kernel =
cu::qmm<DataType, BM, BN, BK, group_size.value, bits.value>;
auto kernel = cu::
qmm_t_aligned<DataType, BM, BN, BK, group_size.value, bits.value>;
dim3 grid(N / BN, M / BM);
enc.add_kernel_node(
kernel,
grid,
128,
2 * 4 * 32,
1 * sizeof(DataType) * (BM * BK + BN * BK),
x.data<DataType>(),
w.data<uint8_t>(),
scales.data<DataType>(),

2
mlx/backend/cuda/random.cu
View File

@@ -170,6 +170,7 @@ void RandomBits::eval_gpu(const std::vector<array>& inputs, array& out) {
cu::rbitsc,
grid,
block,
0,
keys.data<uint32_t>(),
out.data<uint8_t>(),
grid_dims,
@@ -180,6 +181,7 @@ void RandomBits::eval_gpu(const std::vector<array>& inputs, array& out) {
cu::rbits,
grid,
block,
0,
keys.data<uint32_t>(),
out.data<uint8_t>(),
grid_dims,

2
mlx/backend/cuda/reduce/all_reduce.cu
View File

@@ -120,6 +120,7 @@ void all_reduce(
kernel,
blocks,
threads,
0,
static_cast<T*>(indata),
intermediate.data<U>(),
block_step,
@@ -146,6 +147,7 @@ void all_reduce(
kernel,
blocks,
threads,
0,
static_cast<T*>(indata),
out.data<U>(),
block_step,

2
mlx/backend/cuda/reduce/col_reduce.cu
View File

@@ -230,7 +230,7 @@ void col_reduce_looped(
auto kernel =
cu::col_reduce_looped<T, U, OP, reduce_ndim(), BM, BN, N_READS>;
encoder.add_kernel_node(
kernel, grid, blocks, indata, out.data<U>(), args);
kernel, grid, blocks, 0, indata, out.data<U>(), args);
});
});
});

3
mlx/backend/cuda/reduce/init_reduce.cu
View File

@@ -41,7 +41,8 @@ void init_reduce(
dim3 grid = get_2d_grid_dims(out.shape(), out.strides());
dim3 block(grid.x < 1024 ? grid.x : 1024, 1, 1);
grid.x = (grid.x + 1023) / 1024;
encoder.add_kernel_node(kernel, grid, block, out.data<U>(), out.size());
encoder.add_kernel_node(
kernel, grid, block, 0, out.data<U>(), out.size());
});
});
}

4
mlx/backend/cuda/reduce/row_reduce.cu
View File

@@ -269,7 +269,7 @@ void row_reduce_simple(
int size = plan.shape.back();
encoder.add_kernel_node(
kernel, grid, block, indata, out.data<U>(), out.size(), size);
kernel, grid, block, 0, indata, out.data<U>(), out.size(), size);
});
});
}
@@ -322,7 +322,7 @@ void row_reduce_looped(
});
encoder.add_kernel_node(
kernel, grid, block, indata, out.data<U>(), out.size(), args);
kernel, grid, block, 0, indata, out.data<U>(), out.size(), args);
});
});
}

2
mlx/backend/cuda/rms_norm.cu
View File

@@ -232,6 +232,7 @@ void RMSNorm::eval_gpu(
kernel,
n_rows,
block_dim(),
0,
x.data<DataType>(),
w.data<DataType>(),
out.data<DataType>(),
@@ -327,6 +328,7 @@ void RMSNormVJP::eval_gpu(
kernel,
n_rows,
block_dim(),
0,
x.data<DataType>(),
w.data<DataType>(),
g.data<DataType>(),

4
mlx/backend/cuda/rope.cu
View File

@@ -325,6 +325,7 @@ void RoPE::eval_gpu(
kernel,
grid,
block,
0,
(donated ? out : in).data<DataType>(),
out.data<DataType>(),
offset.data<int32_t>(),
@@ -341,6 +342,7 @@ void RoPE::eval_gpu(
kernel,
grid,
block,
0,
(donated ? out : in).data<DataType>(),
out.data<DataType>(),
offset.data<int32_t>(),
@@ -360,6 +362,7 @@ void RoPE::eval_gpu(
kernel,
grid,
block,
0,
(donated ? out : in).data<DataType>(),
out.data<DataType>(),
offset.data<int32_t>(),
@@ -381,6 +384,7 @@ void RoPE::eval_gpu(
kernel,
grid,
block,
0,
(donated ? out : in).data<DataType>(),
out.data<DataType>(),
offset.data<int32_t>(),

2
mlx/backend/cuda/scan.cu
View File

@@ -414,6 +414,7 @@ void Scan::eval_gpu(const std::vector<array>& inputs, array& out) {
kernel,
in.data_size() / axis_size,
block_dim,
0,
in.data<T>(),
out.data<U>(),
axis_size);
@@ -443,6 +444,7 @@ void Scan::eval_gpu(const std::vector<array>& inputs, array& out) {
kernel,
num_blocks,
block_dim,
0,
in.data<T>(),
out.data<U>(),
axis_size,

1
mlx/backend/cuda/softmax.cu
View File

@@ -152,6 +152,7 @@ void Softmax::eval_gpu(const std::vector<array>& inputs, array& out) {
kernel,
n_rows,
block_dim(),
0,
in.data<DataType>(),
out.data<DataType>(),
axis_size);

3
mlx/backend/cuda/ternary.cu
View File

@@ -133,6 +133,7 @@ void ternary_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
a.data<bool>(),
b.data<DType>(),
c.data<DType>(),
@@ -151,6 +152,7 @@ void ternary_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
a.data<bool>(),
b.data<DType>(),
c.data<DType>(),
@@ -180,6 +182,7 @@ void ternary_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
a.data<bool>(),
b.data<DType>(),
c.data<DType>(),

2
mlx/backend/cuda/unary.cu
View File

@@ -142,6 +142,7 @@ void unary_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
in.data<InType>(),
out.data<OutType>(),
out.data_size());
@@ -154,6 +155,7 @@ void unary_op_gpu_inplace(
kernel,
num_blocks,
block_dims,
0,
in.data<InType>(),
out.data<OutType>(),
out.data_size(),