// 明确指定线程组织dim3 gridSize(blocksX, blocksY);dim3 blockSize(threadsPerBlockX, threadsPerBlockY);matmul_kernel<<<gridSize, blockSize>>>(A, B, C, M, N, K);

tile→描述数据分块

# cuTile Python - 只需描述数据分块@cutile.primitivedef matmul(M: cutile.dim, N: cutile.dim, K: cutile.dim,           A: cutile.tensor[M, K], B: cutile.tensor[K, N], C: cutile.tensor[M, N]):    # 定义Tile大小    tile_m = cutile.constant(128)    tile_n = cutile.constant(128)    tile_k = cutile.constant(32)    # 自动生成分块计算    for i in cutile.axis(0, M, tile_m):      # 外层循环自动并行化        for j in cutile.axis(0, N, tile_n):  # 编译器决定并行策略            accum = cutile.allocate((tile_m, tile_n), "float32")            for k in cutile.axis(0, K, tile_k):                a_tile = A[i:i+tile_m, k:k+tile_k]                b_tile = B[k:k+tile_k, j:j+tile_n]                accum = cutile.dot(a_tile, b_tile, accum)            C[i:i+tile_m, j:j+tile_n] = accum

2.内存访问模式

传统cuda→→手动、手动，真麻烦

// 需要显式管理共享内存__shared__ float tileA[TILE_SIZE][TILE_SIZE];__shared__ float tileB[TILE_SIZE][TILE_SIZE];// 手动加载到共享内存int tx = threadIdx.x, ty = threadIdx.y;tileA[ty][tx] = A[row * K + col];tileB[ty][tx] = B[row * N + col];__syncthreads();  // 必须显式同步

tile→→→自动、自动→全自动

# 编译器自动处理内存层次@cutile.primitivedef conv2d(input: cutile.tensor[N, H, W, C],            filter: cutile.tensor[Fh, Fw, C, M],           output: cutile.tensor[N, H_out, W_out, M]):    # 编译器决定何时使用共享内存、寄存器    tile_n = 8   # batch维度分块    tile_h = 16  # 高度分块    tile_w = 16  # 宽度分块    tile_c = 32  # 通道分块    # 自动生成内存层次代码    for n, h, w, m in cutile.grid(N, H_out, W_out, M):        for fh, fw, c in cutile.reduction(Fh, Fw, C):            # 编译器决定内存放置策略            in_tile = input[n, h*stride+fh, w*stride+fw, c]            filter_tile = filter[fh, fw, c, m]            output[n, h, w, m] += in_tile * filter_tile

3.并行循环表达

传统cuda

// 网格/线程块映射int row = blockIdx.y * blockDim.y + threadIdx.y;int col = blockIdx.x * blockDim.x + threadIdx.x;// 并行循环需要手动展开if (row < M && col < N) {    float sum = 0.0f;    for (int k = 0; k < K; k++) {        sum += A[row * K + k] * B[k * N + col];    }    C[row * N + col] = sum;}

tile（瓦片）

# 声明并行维度@cutile.primitivedef batched_gemm(batch: cutile.dim, M: cutile.dim, N: cutile.dim, K: cutile.dim,                 A: cutile.tensor[batch, M, K], B: cutile.tensor[batch, K, N],                 C: cutile.tensor[batch, M, N]):    # cutile.axis 定义并行维度    for b in cutile.axis(0, batch, 8):  # batch维度并行        for i in cutile.axis(0, M, 128):  # M维度并行            for j in cutile.axis(0, N, 128):  # N维度并行                # reduction维度自动处理                accum = cutile.sum_reduce(K, 32,                     lambda k: A[b, i, k] * B[b, k, j])                C[b, i, j] = accum