行指针的使用——*(*(pointer + row) + col)是什么意思

最新推荐文章于 2023-08-18 17:23:34 发布
原创 最新推荐文章于 2023-08-18 17:23:34 发布 · 1.4k 阅读 · 3 ·
CC 4.0 BY-SA版权
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
文章标签：

#c #c++ #行指针

本文探讨了如何使用行指针来访问二维数组中的元素，特别是解释了表达式`*(*(pointer + row) + col)`的含义。通过一个实例，解释了如何用行指针对3X3二维数组的每一行找到最大值并存储到一维数组中。

行指针的使用——((pointer + row) + col)是什么意思
用指针分别指向二维数组中的一个元素和二维数组中的一个子数组的时候，可以形象的称指向单个元素的为列指针，指向子数组的指针为行指针。

#include<stdio.h>

int main()
{
   
   
	int array[3][4] = {
   
    {
   
   1,2,3,4},{
   
   5,6,7,8},{
   
   9,10,11,12} };
	int<
最低0.47元/天 解锁文章
Karon_R
关注
行指针的奇葩用法
Lontern的博客
09-22 443
行指针的奇葩用法
行指针（数组指针）及其运用
weixin_43399456的博客
04-16 576
方括号中的常量表达式表示指针加1，移动几个数据；即指针加一，移动三个数据。当用行指针操作二维数组时，表达式一般写成1行的元素个数，即列数。用行指针表示二维数组的某个元素，并且遍历。存储行地址的指针变量，叫做行指针变量。例如 int (*p)[3];
行指针
hongqianshen的博客
09-14 5164
行指针一般用于二维数组中，指向每行元素的首地址； 定义形式：int (*p) [5] ; 表示方法：直接在p上偏移就是行，在*p上偏移就是列 #include void main(){     int n[][3] = {10,20,30,40,50,60};     int (*p)[3];     p=n;     cout",""," }  
行指针和列指针-专题分析(转至汪老师的C语言)
haussuden的专栏
10-25 2万+
<br />指针与多维数组<br />（主要指二维数组）<br />int a[3][4]={1,3,5,7,9,11,13,15,17,19,21,23};<br /> <br /><br />换个角度看世界：<br />如首行一样，将首行视为一个元素，一个特殊的元素，这个“特殊的”元素是一个一维数组。那么这个二维数组是由是由三个“特殊的”元素组成的一个“特殊的”一维数组。<br />a是这个“特殊的”一维数组的名称，也就是首地址，也就是第一个元素的地址，也就是第一行的首地址，是指首行一整行，并不是指某个
指针用法及其详解
云龙的博客
02-10 3166
指针与数组 我们举一个变相引用指针的例子：数组名是数组首元素的地址。 例如：假如a是一个数组，则下面的语句成立 a = &a[0];//数组名是数组首元素 的地址 下面的等式说说明了数组和指针的关系： data + 2 == &data[2] //相同的地址 *(data+2) ==data[2] //相同的值 C语言在描述数组的时候借助了指针，也就是说，定义a[n]的...
IntPtr imgPtr = imgMemory.Scan0 + (i + j) * imgMemory.Stride + x1;
最新发布
12-06
而原代码中的(i+j)相当于这里的row（总行偏移），x1相当于col*pixelSize（列字节偏移）。 12. **与问题中代码的对应**： 原代码：IntPtr imgPtr = imgMemory.Scan0 + (i + j) * imgMemory.Stride + x1; 这里x1...
S32 refine_process_char(REFINE_CONTEXT *refine_config, FG *src_image, FG *dst_image, float stop_percent) { if ((refine_config == NULL) || (src_image == NULL) || (dst_image == NULL)) { AMS_ERROR("refine process Input is NULL"); return ERROR; } S32 sum = 0; U16 width = refine_config->width; U16 height = refine_config->height; U8 radius = refine_config->radius; S32 stop_amount = (2 * radius + 1) * (2 * radius + 1) * stop_percent; U16 *col_hist = refine_config->col_hist; U16 *hist = refine_config->hist; S32 *row_offset = refine_config->row_offset; S32 *col_offset = refine_config->col_offset; get_offset_pos(row_offset, width, radius, radius); get_offset_pos(col_offset, height, radius, radius); memset(col_hist, 0, 2 * (width + radius + radius) * sizeof(U16)); S16 row, col; for (row = 0; row < height; row++) { if (row == 0) { for (S16 k = -radius; k <= radius; k++) { U8 *line_ps = src_image + col_offset[k + radius] * width; for (col = -radius; col < width + radius; col++) { col_hist[(col + radius) * 2 + line_ps[row_offset[col + radius]]]++; } } } else { U8 *line_ps = src_image + col_offset[row - 1] * width; for (col = -radius; col < width + radius; col++) { col_hist[(col + radius) * 2 + line_ps[row_offset[col + radius]]]--; } line_ps = src_image + col_offset[row + radius + radius] * width; for (col = -radius; col < width + radius; col++) { col_hist[(col + radius) * 2 + line_ps[row_offset[col + radius]]]++; } } memset(hist, 0, 2 * sizeof(uint16_t)); U8 *line_pd = dst_image + row * width; for (col = 0; col < width; col++) { if (col == 0) { for (S16 k = -radius; k <= radius; k++) { U16 *col_ptr = col_hist + (k + radius) * 2; hist[0] = hist[0] + col_ptr[0]; hist[1] = hist[1] + col_ptr[1]; } } else { U16 *col_ptr1 = col_hist + (row_offset[col - 1] + radius) * 2; U16 *col_ptr2 = col_hist + (row_offset[col + radius + radius] + radius) * 2; hist[0] = col_ptr2[0] + hist[0] - col_ptr1[0]; hist[1] = col_ptr2[1] + hist[1] - col_ptr1[1]; } sum = 0; for (U16 i = 0; i < 2; i++) { sum += hist[i]; if (sum >= stop_amount) { *line_pd = i; break; } } line_pd++; } } return OK; } /* Pixel mirroring */ LOCAL S32 get_mirror_pos(U16 length, S32 pos) { if (pos < 0) { pos = -pos; while (pos > length) { pos -= length; } } else if (pos >= length) { pos = length - (pos - length + 2); while (pos < 0) { pos += length; } } return pos; } LOCAL void get_offset_pos(S32 *pos_ptr, U16 length, U16 left, U16 right) { for (S32 col = -left; col < length + right; col++) { pos_ptr[col + left] = get_mirror_pos(length, col); } } 这段代码实现了在全图内进行refine，将其改造为在一个窗口内（left right top bottom）进行refine的代码
10-12
在else部分，它减去row - 1（上边界），添加row + 2*radius（下边界），因为对于当前行row，kernel的行范围是[row - radius, row + radius]，所以当向下移动一行时，移除row - radius - 1? 错误。 假设kernel大小为...
import torch import triton import triton.language as tl @triton.jit def fused_bmm_softmax_kernel( # 输入/输出指针 a_ptr, b_ptr, c_ptr, # 张量维度 B, M, N, K, # 张量A步长 (batch, row, col) stride_ab, stride_am, stride_ak, # 张量B步长 (batch, col, col) stride_bb, stride_bk, stride_bn, # 输出C步长 (batch, row, col) stride_cb, stride_cm, stride_cn, # 分块参数 BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr ): # 程序ID对应输出矩阵中的位置 (batch, row) batch_id = tl.program_id(0) row_id = tl.program_id(1) # 初始化行最大值为负无穷 row_max = tl.full((), float('-inf'), dtype=tl.float32) # Pass 1: 计算当前行的最大值 for n_block_idx in range(0, tl.cdiv(N, BLOCK_SIZE_N)): # 为当前列块预加载A的部分数据 a_ptrs = a_ptr + batch_id * stride_ab + row_id * stride_am + \ tl.arange(0, BLOCK_SIZE_K)[None, :] * stride_ak b_ptrs = b_ptr + batch_id * stride_bb + \ tl.arange(0, BLOCK_SIZE_K)[:, None] * stride_bk + \ (n_block_idx * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)[None, :]) * stride_bn # 累加部分矩阵乘法结果 accumulator = tl.zeros((BLOCK_SIZE_N,), dtype=tl.float32) for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)): a = tl.load(a_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[None, :] < K - k * BLOCK_SIZE_K, other=0.0) b = tl.load(b_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[:, None] < K - k * BLOCK_SIZE_K, other=0.0) partial = tl.dot(a, b) accumulator += partial a_ptrs += BLOCK_SIZE_K * stride_ak b_ptrs += BLOCK_SIZE_K * stride_bk # 更新行最大值 max_val = tl.max(accumulator) row_max = tl.maximum(row_max, max_val) # Pass 2: 计算指数和 (exp_sum) exp_sum = tl.zeros((), dtype=tl.float32) for n_block_idx in range(0, tl.cdiv(N, BLOCK_SIZE_N)): # 重新加载数据（与Pass 1类似） a_ptrs = a_ptr + batch_id * stride_ab + row_id * stride_am + \ tl.arange(0, BLOCK_SIZE_K)[None, :] * stride_ak b_ptrs = b_ptr + batch_id * stride_bb + \ tl.arange(0, BLOCK_SIZE_K)[:, None] * stride_bk + \ (n_block_idx * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)[None, :]) * stride_bn accumulator = tl.zeros((BLOCK_SIZE_N,), dtype=tl.float32) for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)): a = tl.load(a_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[None, :] < K - k * BLOCK_SIZE_K, other=0.0) b = tl.load(b_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[:, None] < K - k * BLOCK_SIZE_K, other=0.0) partial = tl.dot(a, b) accumulator += partial a_ptrs += BLOCK_SIZE_K * stride_ak b_ptrs += BLOCK_SIZE_K * stride_bk # 计算稳定指数值并累加 exp_vals = tl.exp(accumulator - row_max) exp_sum += tl.sum(exp_vals) # Pass 3: 计算归一化结果并写入 for n_block_idx in range(0, tl.cdiv(N, BLOCK_SIZE_N)): a_ptrs = a_ptr + batch_id * stride_ab + row_id * stride_am + \ tl.arange(0, BLOCK_SIZE_K)[None, :] * stride_ak b_ptrs = b_ptr + batch_id * stride_bb + \ tl.arange(0, BLOCK_SIZE_K)[:, None] * stride_bk + \ (n_block_idx * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)[None, :]) * stride_bn accumulator = tl.zeros((BLOCK_SIZE_N,), dtype=tl.float32) for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)): a = tl.load(a_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[None, :] < K - k * BLOCK_SIZE_K, other=0.0) b = tl.load(b_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[:, None] < K - k * BLOCK_SIZE_K, other=0.0) partial = tl.dot(a, b) accumulator += partial a_ptrs += BLOCK_SIZE_K * stride_ak b_ptrs += BLOCK_SIZE_K * stride_bk # 计算最终Softmax结果 exp_vals = tl.exp(accumulator - row_max) softmax_output = exp_vals / exp_sum # 写入输出 c_ptrs = c_ptr + batch_id * stride_cb + row_id * stride_cm + \ (n_block_idx * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) * stride_cn tl.store(c_ptrs, softmax_output) 封装函数供PyTorch调用 def fused_bmm_softmax(a: torch.Tensor, b: torch.Tensor): B, M, K = a.shape B, K, N = b.shape c = torch.empty((B, M, N), device=a.device, dtype=a.dtype) # 配置分块大小 (需根据硬件调整) BLOCK_SIZE_M = 1 # 每个程序实例处理1行 BLOCK_SIZE_N = 64 # 列方向分块大小 BLOCK_SIZE_K = 32 # K维度分块大小 grid = (B, M) # 每个批次每行一个程序实例 fused_bmm_softmax_kernel[grid]( a, b, c, B, M, N, K, a.stride(0), a.stride(1), a.stride(2), b.stride(0), b.stride(1), b.stride(2), c.stride(0), c.stride(1), c.stride(2), BLOCK_SIZE_M=BLOCK_SIZE_M, BLOCK_SIZE_N=BLOCK_SIZE_N, BLOCK_SIZE_K=BLOCK_SIZE_K ) return c把上面的融合kernel跟据下面的代码的思路写一个配套的代码（完全与融合kernel配套，只是代码思路类似而已）：import os import sys import random import itertools import csv from triton import Config import torch import triton import torch_mlu import torch.nn as nn import triton.language as tl os.environ[“TRITON_PRINT_AUTOTUNING”] = “1” def get_config(): row_block = list(range(1, 33)) num_stages = [1, 2, 3, 4] configs = list(itertools.product(row_block, num_stages)) random.shuffle(configs) return configs @triton.jit def softmax_kernel(output_ptr, input_ptr, input_row_stride, output_row_stride, n_rows, n_cols, BLOCK_SIZE: tl.constexpr, OUTER_ROW_BLOCK: tl.constexpr, ROW_BLOCK: tl.constexpr, num_stages: tl.constexpr): # The rows of the softmax are independent, so we parallelize across rows. pid = tl.program_id(0) ub = tl.minimum(n_rows, (tl.program_id(0) + 1) * OUTER_ROW_BLOCK) outer_raw_idx = tl.program_id(0) * OUTER_ROW_BLOCK for inner_row in range(outer_raw_idx, ub, ROW_BLOCK): row_offsets = tl.arange(0, ROW_BLOCK) row_mask = (inner_row + row_offsets < ub)[:, None] # The stride represents how much we need to increase the pointer to advance 1 row row_start_ptr = input_ptr + (inner_row + row_offsets) * input_row_stride col_offsets = tl.arange(0, BLOCK_SIZE)[None, :] input_ptrs = row_start_ptr[:, None] + col_offsets # Load rows into SRAM row = tl.load(input_ptrs, mask=row_mask).to(tl.float32) # Subtract maximum for numerical stability row_minus_max = row - tl.max(row, axis=1)[:, None] # Note that exponentiation in Triton is fast but approximate numerator = tl.exp(row_minus_max) denominator = tl.sum(numerator, axis=1)[:, None] softmax_output = numerator / denominator softmax_output = softmax_output.to(tl.float16) # Write back output to DRAM output_row_start_ptr = output_ptr + (inner_row + row_offsets) * output_row_stride output_ptrs = output_row_start_ptr[:, None] + col_offsets tl.store(output_ptrs, softmax_output, mask=row_mask) def triton_softmax_perf(configs, x): n_rows, n_cols = x.shape BLOCK_SIZE = n_cols # Allocate output y = torch.empty_like(x) # Enqueue kernel. We split all rows into 48 parts evenly. OUTER_ROW_BLOCK = triton.cdiv(n_rows, 48) grid = lambda META: (48, 1, 1) softmax_kernel[(grid)](y, x, x.stride(0), y.stride(0), n_rows, n_cols, BLOCK_SIZE=BLOCK_SIZE, OUTER_ROW_BLOCK=OUTER_ROW_BLOCK, ROW_BLOCK=configs[0], num_stages=configs[1]) return y def benchmark_softmax(configs, M, K, warmup, repeat): # 创建输入数据 a = torch.randn(M, K, dtype=torch.float16, device=‘mlu’).contiguous() trainset, count, maxsampled = [], 0, 20 for i in range(len(configs)): try: # filter invalid settings c_triton = triton_softmax_perf(configs[i], a) except: continue for _ in range(warmup): c_triton = triton_softmax_perf(configs[i], a) torch.mlu.synchronize() start_event = torch.mlu.Event(enable_timing=True) end_event = torch.mlu.Event(enable_timing=True) start_event.record() for _ in range(repeat): c_triton = triton_softmax_perf(configs[i], a) torch.mlu.synchronize() end_event.record() torch.mlu.synchronize() trainset.append(list(configs[i])) trainset[count].append(start_event.hardware_time(end_event) / repeat / 1000) print('train data {}: {}'.format(count, trainset[count])) count += 1 if count >= maxsampled: break with open('data/{}_{}.csv'.format(M, K), 'w', newline='') as f: writer = csv.writer(f) writer.writerows(trainset) def main(): micro_shapes = [ (80, 768), (604, 768), (1216, 768), (1968, 768), ] configs = get_config() for M, K in micro_shapes: print('Processing input {} {}'.format(M, K)) benchmark_softmax(configs, M, K, warmup=10, repeat=20) if name==“main”: main()
06-12
# 张量A步长 (batch, row, col) stride_ab, stride_am, stride_ak, # 张量B步长 (batch, col, col) stride_bb, stride_bk, stride_bn, # 输出C步长 (batch, row, col) stride_cb, stride_cm, stride_cn, # ...
import torch import triton import triton.language as tl @triton.jit def fused_bmm_softmax_kernel( # 输入/输出指针 a_ptr, b_ptr, c_ptr, # 张量维度 B, M, N, K, # 张量A步长 (batch, row, col) stride_ab, stride_am, stride_ak, # 张量B步长 (batch, col, col) stride_bb, stride_bk, stride_bn, # 输出C步长 (batch, row, col) stride_cb, stride_cm, stride_cn, # 分块参数 BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr ): # 程序ID对应输出矩阵中的位置 (batch, row) batch_id = tl.program_id(0) row_id = tl.program_id(1) # 初始化行最大值为负无穷 row_max = tl.full((), float('-inf'), dtype=tl.float32) # Pass 1: 计算当前行的最大值 for n_block_idx in range(0, tl.cdiv(N, BLOCK_SIZE_N)): # 为当前列块预加载A的部分数据 a_ptrs = a_ptr + batch_id * stride_ab + row_id * stride_am + \ tl.arange(0, BLOCK_SIZE_K)[None, :] * stride_ak b_ptrs = b_ptr + batch_id * stride_bb + \ tl.arange(0, BLOCK_SIZE_K)[:, None] * stride_bk + \ (n_block_idx * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)[None, :]) * stride_bn # 累加部分矩阵乘法结果 accumulator = tl.zeros((BLOCK_SIZE_N,), dtype=tl.float32) for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)): a = tl.load(a_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[None, :] < K - k * BLOCK_SIZE_K, other=0.0) b = tl.load(b_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[:, None] < K - k * BLOCK_SIZE_K, other=0.0) partial = tl.dot(a, b) accumulator += partial a_ptrs += BLOCK_SIZE_K * stride_ak b_ptrs += BLOCK_SIZE_K * stride_bk # 更新行最大值 max_val = tl.max(accumulator) row_max = tl.maximum(row_max, max_val) # Pass 2: 计算指数和 (exp_sum) exp_sum = tl.zeros((), dtype=tl.float32) for n_block_idx in range(0, tl.cdiv(N, BLOCK_SIZE_N)): # 重新加载数据（与Pass 1类似） a_ptrs = a_ptr + batch_id * stride_ab + row_id * stride_am + \ tl.arange(0, BLOCK_SIZE_K)[None, :] * stride_ak b_ptrs = b_ptr + batch_id * stride_bb + \ tl.arange(0, BLOCK_SIZE_K)[:, None] * stride_bk + \ (n_block_idx * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)[None, :]) * stride_bn accumulator = tl.zeros((BLOCK_SIZE_N,), dtype=tl.float32) for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)): a = tl.load(a_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[None, :] < K - k * BLOCK_SIZE_K, other=0.0) b = tl.load(b_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[:, None] < K - k * BLOCK_SIZE_K, other=0.0) partial = tl.dot(a, b) accumulator += partial a_ptrs += BLOCK_SIZE_K * stride_ak b_ptrs += BLOCK_SIZE_K * stride_bk # 计算稳定指数值并累加 exp_vals = tl.exp(accumulator - row_max) exp_sum += tl.sum(exp_vals) # Pass 3: 计算归一化结果并写入 for n_block_idx in range(0, tl.cdiv(N, BLOCK_SIZE_N)): a_ptrs = a_ptr + batch_id * stride_ab + row_id * stride_am + \ tl.arange(0, BLOCK_SIZE_K)[None, :] * stride_ak b_ptrs = b_ptr + batch_id * stride_bb + \ tl.arange(0, BLOCK_SIZE_K)[:, None] * stride_bk + \ (n_block_idx * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)[None, :]) * stride_bn accumulator = tl.zeros((BLOCK_SIZE_N,), dtype=tl.float32) for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)): a = tl.load(a_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[None, :] < K - k * BLOCK_SIZE_K, other=0.0) b = tl.load(b_ptrs, mask=tl.arange(0, BLOCK_SIZE_K)[:, None] < K - k * BLOCK_SIZE_K, other=0.0) partial = tl.dot(a, b) accumulator += partial a_ptrs += BLOCK_SIZE_K * stride_ak b_ptrs += BLOCK_SIZE_K * stride_bk # 计算最终Softmax结果 exp_vals = tl.exp(accumulator - row_max) softmax_output = exp_vals / exp_sum # 写入输出 c_ptrs = c_ptr + batch_id * stride_cb + row_id * stride_cm + \ (n_block_idx * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) * stride_cn tl.store(c_ptrs, softmax_output) # 封装函数供PyTorch调用 def fused_bmm_softmax(a: torch.Tensor, b: torch.Tensor): B, M, K = a.shape B, K, N = b.shape c = torch.empty((B, M, N), device=a.device, dtype=a.dtype) # 配置分块大小 (需根据硬件调整) BLOCK_SIZE_M = 1 # 每个程序实例处理1行 BLOCK_SIZE_N = 64 # 列方向分块大小 BLOCK_SIZE_K = 32 # K维度分块大小 grid = (B, M) # 每个批次每行一个程序实例 fused_bmm_softmax_kernel[grid]( a, b, c, B, M, N, K, a.stride(0), a.stride(1), a.stride(2), b.stride(0), b.stride(1), b.stride(2), c.stride(0), c.stride(1), c.stride(2), BLOCK_SIZE_M=BLOCK_SIZE_M, BLOCK_SIZE_N=BLOCK_SIZE_N, BLOCK_SIZE_K=BLOCK_SIZE_K ) return c把上面的融合kernel替换到下面的代码中：import os import sys import random import itertools import csv from triton import Config import torch import triton import torch_mlu import torch.nn as nn import triton.language as tl os.environ["TRITON_PRINT_AUTOTUNING"] = "1" def get_config(): row_block = list(range(1, 33)) num_stages = [1, 2, 3, 4] configs = list(itertools.product(row_block, num_stages)) random.shuffle(configs) return configs @triton.jit def softmax_kernel(output_ptr, input_ptr, input_row_stride, output_row_stride, n_rows, n_cols, BLOCK_SIZE: tl.constexpr, OUTER_ROW_BLOCK: tl.constexpr, ROW_BLOCK: tl.constexpr, num_stages: tl.constexpr): # The rows of the softmax are independent, so we parallelize across rows. pid = tl.program_id(0) ub = tl.minimum(n_rows, (tl.program_id(0) + 1) * OUTER_ROW_BLOCK) outer_raw_idx = tl.program_id(0) * OUTER_ROW_BLOCK for inner_row in range(outer_raw_idx, ub, ROW_BLOCK): row_offsets = tl.arange(0, ROW_BLOCK) row_mask = (inner_row + row_offsets < ub)[:, None] # The stride represents how much we need to increase the pointer to advance 1 row row_start_ptr = input_ptr + (inner_row + row_offsets) * input_row_stride col_offsets = tl.arange(0, BLOCK_SIZE)[None, :] input_ptrs = row_start_ptr[:, None] + col_offsets # Load rows into SRAM row = tl.load(input_ptrs, mask=row_mask).to(tl.float32) # Subtract maximum for numerical stability row_minus_max = row - tl.max(row, axis=1)[:, None] # Note that exponentiation in Triton is fast but approximate numerator = tl.exp(row_minus_max) denominator = tl.sum(numerator, axis=1)[:, None] softmax_output = numerator / denominator softmax_output = softmax_output.to(tl.float16) # Write back output to DRAM output_row_start_ptr = output_ptr + (inner_row + row_offsets) * output_row_stride output_ptrs = output_row_start_ptr[:, None] + col_offsets tl.store(output_ptrs, softmax_output, mask=row_mask) def triton_softmax_perf(configs, x): n_rows, n_cols = x.shape BLOCK_SIZE = n_cols # Allocate output y = torch.empty_like(x) # Enqueue kernel. We split all rows into 48 parts evenly. OUTER_ROW_BLOCK = triton.cdiv(n_rows, 48) grid = lambda META: (48, 1, 1) softmax_kernel[(grid)](y, x, x.stride(0), y.stride(0), n_rows, n_cols, BLOCK_SIZE=BLOCK_SIZE, OUTER_ROW_BLOCK=OUTER_ROW_BLOCK, ROW_BLOCK=configs[0], num_stages=configs[1]) return y def benchmark_softmax(configs, M, K, warmup, repeat): # 创建输入数据 a = torch.randn(M, K, dtype=torch.float16, device='mlu').contiguous() trainset, count, maxsampled = [], 0, 20 for i in range(len(configs)): try: # filter invalid settings c_triton = triton_softmax_perf(configs[i], a) except: continue for _ in range(warmup): c_triton = triton_softmax_perf(configs[i], a) torch.mlu.synchronize() start_event = torch.mlu.Event(enable_timing=True) end_event = torch.mlu.Event(enable_timing=True) start_event.record() for _ in range(repeat): c_triton = triton_softmax_perf(configs[i], a) torch.mlu.synchronize() end_event.record() torch.mlu.synchronize() trainset.append(list(configs[i])) trainset[count].append(start_event.hardware_time(end_event) / repeat / 1000) print('train data {}: {}'.format(count, trainset[count])) count += 1 if count >= maxsampled: break with open('data/{}_{}.csv'.format(M, K), 'w', newline='') as f: writer = csv.writer(f) writer.writerows(trainset) def main(): micro_shapes = [ (80, 768), (604, 768), (1216, 768), (1968, 768), ] configs = get_config() for M, K in micro_shapes: print('Processing input {} {}'.format(M, K)) benchmark_softmax(configs, M, K, warmup=10, repeat=20) if __name__=="__main__": main()
06-12
在softmax中，我们通常是对行做softmax，所以矩阵乘法的结果应该是(1,M,N)，然后对每一行做softmax。在我们的融合kernel中，N可以任意指定。为了和原始基准测试一致，我们可以让N等于M（即输出为方阵）？但原始基准...
C语言之指针用法
努力跟上大佬脚步的小码农酥酥
08-18 574
指针的用法
三大指针：行指针、指针数组、二级指针
百慕大
06-07 1万+
“三大指针”这个是自己命名的，应该是指针中比较难理解的一块，所以拿出来写了个笔记。   友情提示：在看下面部分之前，需要对指针指向二维数组有个比较清楚的了解   在介绍三个指针之前，一定要牢记下面这个图，也就是二级指针的概念。 1、行指针：数据类型(*p)[m]; 定义：指向由m个元素组成的一维数组的行指针变量 经测试，不能直接指向一维数组，如： int a[3];
C语言的行指针
sevanco的博客
03-26 2294
今天复习指针的时候看到书上讲二维数组中的两类指针： 一类是元素指针 一类是行指针 元素指针指向一个元素，行指针指向整个行。（这里比较难以理解的是行指针）行指针可以转化为元素指针，方法是在前面加一个”∗*”,指向改行的第一个元素。也就是说加上了”∗*”之后，它还是一个地址。代码#include <stdio.h> #include <string.h> #define M 2 #define N 2
数组指针 行指针 列指针
笔记
02-26 3220
概念 我们把指向数组的指针叫做数组指针，后面还会学到指针数组，这两个是不一样的，根据中学语文偏正词组的知识可以知道，前者是指针，后者是数组。 一般指针变量 int a[2][3]={{1,2,3},{4,5,6}}; int *P=&a[0][0];//int *p=a[0]; 数组在内存中的存储都是连续的，并且是先存放好第一行的元素，再放第二行的元素，因此，p[0]和*p就是a[0][0],p[5]和*(p+5)就是a[1][2]。 int a[m][n];即，a[i][j]可用*(p
C语言：行指针和列指针
热门推荐
sudazf的专栏
03-11 4万+
int* a; //a是指向整形的指针； int* a[5]; //a是一个一维数组名（该数组包含5个元素，每个元素类型是指向整形的指针）； //a表示的是数组（a称指针数组）； //从数组元素角度看，a的值是首元素地址，a+1是第二个元素地址...； int (*a)[5]; //a是指向一维数组的指针（这里每个一维数组含5个整形元素）； //a的值是第一个一维数组的首元素地址,a+1是第二个一维数组的首元素地址...（a是数组指针）；
行指针和列指针的理解
Jack
06-05 2万+
指针与二维数组 int a[3][4]={1,3,5,7,9,11,13,15,17,19,21,23}; 我们以为，二维数组的存储形式是： 但是，二维数组实际上是以一维数组为单位连续存储的。 这样，上述的二维数组就可以看成是“特殊的”一维数组了。 而数组名a是这个“特殊的”一维数组的名称，也就是首地址（常量地址），也就是第一个元素的地址。 也就是第一行的首地址，是指首行一整行，...
【C++】二维数组的行指针和列指针
zzy979481894的博客
05-11 1万+
在C++中，数组与指针有密切的关系。任何通过数组和下标实现的表达式可等价地通过指针和偏移量实现。下面首先介绍一维数组与指针的关系，之后介绍二维数组的行指针和列指针。
指针的使用总结
编程猿来如此
07-10 728
指针的使用总结 指针的概念 内存地址 内存是一个线性的字节数组，一个字节由 8 个二进制位组成，每个字节都有唯一的编号，这个编号就是内存地址，代表一个字节的的存储空间。 指针的概念 指针是“指向”另外一种类型的复合类型，可以实现对其他对象的间接访问，指针的本质就是数据在内存中的地址，地址即指针。 指针变量 概念 指针变量就是用来保存对象地址的变量。如果指针 ptr 保存了 num 的地址，就是说 ptr 指向了 num，也就是 ptr 指向了 num 所在的内存块。 定义指针变量 定义指针类
c语言行指针如何表示,关于C语言行指针的详细解析
weixin_34063855的博客
05-19 3698
From:http://soft.chinabyte.com/database/339/12423339.shtml最近在百度知道里面发现很多学习初学者搞不清楚行指针的具体概念，今天在这里作下详细的解释分析。行指针，顾名思义就是指向一行的指针。那么哪里会用到行指针呢，用的最多的，那就是二维指针了，大家都知道，我们通常把二维指针看成一个行列式，但是它在内存中的排序却是和一维指针一样的。比如组a[2]...
指针小结
逆时针
03-11 407
目录 目录 声明指针 指针赋值 给指针解除引用 指针和指针指向的值 指针和数组 指针算术 指针和字符串 链表使用 C++管理数据内存的方式 二维指针 本篇文章记录一些指针学习的心得总结。 声明指针 //typeName *pointerName //for example: int * pointer1; double *pointer2; 指针赋值...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值