cuda ptx 汇编语言示例：读寄存器

编译 , Ampere 显卡，rtx 3060 3070...

nvcc -arch=sm_86 -o hello hello_ptx.cu

或写成Makefile：

hello: hello_sm_id.cunvcc -arch=sm_86 -o $@ $^
#nvcc -arch=sm_86 -o hello hello_sm_id.cu

$@ 是指目标

$^ 是指第一个依赖 ^^

hello_ptx.cu

#include <stdio.h>
#include <stdint.h>static __device__ __inline__ uint32_t __mysmid(){uint32_t ssmid;asm volatile("mov.u32 %0, %%smid;" : "=r"(ssmid));return ssmid;}static __device__ __inline__ uint32_t __mywarpid(){uint32_t warpid;asm volatile("mov.u32 %0, %%warpid;" : "=r"(warpid));return warpid;}static __device__ __inline__ uint32_t __mylaneid(){uint32_t laneid;asm volatile("mov.u32 %0, %%laneid;" : "=r"(laneid));return laneid;}__global__ void mykernel(){int idx = threadIdx.x+blockDim.x*blockIdx.x;unsigned thx = threadIdx.x;
//      if(threadIdx.x==1023)// && blockIdx.x<3)
//      if(threadIdx.x==1)
//      if((thx==0 || thx==32 || thx==64 || thx==96 || thx==128 || thx==160 || thx==192 || thx==224 || thx==256) && blockIdx.x==0)if(blockIdx.x<=33 && thx==0)printf("I am thread %d, my SM ID is %d, my warp ID is %d, and my warp lane is %d\n", idx, __mysmid(), __mywarpid(), __mylaneid());
}int main(){dim3 grid_;dim3 block_;grid_.x=34;block_.x=1024;mykernel<<<grid_,block_>>>();cudaDeviceSynchronize();return 0;
}
//$ nvcc -arch=sm_20 -o hello hello_ptx.cu

运行：

结果分析：

rtx3060中，

当 block 的个数从一个增加到两个，3个，...

smid的值为 0, 2, 4， ...直到偶数最大 max_even(smid), 然后是 1，3，5，... 直到奇数最大 max_odd(smid).

rtx 3060 2 SM/TPC

奇偶的变化，跟这个里的2是否有关系呢？以及有什么关系呢？

cuda ptx 汇编语言示例：读寄存器相关推荐

【Android 逆向】代码调试器开发 ( ptrace 函数 | 读寄存器 | 写寄存器 )
文章目录一.读寄存器二.写寄存器一.读寄存器调用 ptrace(PTRACE_GETREGS, m_nPid, NULL, regs) 读取进程运行时的寄存器 ; 读取寄存器时 , 进程必须处 ...
STM32开发笔记112：ADS1258驱动设计——读寄存器
文章目录前言 1.转换速率 2.寄存器读命令 (1)时序 (2)命令字节 (3)读1字节驱动程序 (4)读寄存器驱动程序 (5)测试程序 (6)上电缺省值 3.寄存器详解 (1)CONFIG0 BI ...
CUDA PTX ISA阅读笔记（一）
不知道这是个啥的看这里:Parallel Thread Execution ISA Version 5.0. 简要来说,PTX就是.cu代码编译出来的一种东西,然后再由PTX编译生成执行代码.如果不想 ...
寄存器和存储器的区别_汇编语言第二章寄存器
第二章寄存器在 CPU 中: 运算器进行信息处理: 寄存器进行信息存储(主要部分,工作原理): 控制器控制各种器件进行工作: 内部总线连接各种器件,在它们之间进行数据的传送. 不同的 CPU,寄存 ...
c++ nvcc编译CUDA程序入门示例
nvcc nvcc是NVIDIA CUDA Compiler,用来编译host和device程序. 这里的术语: host:指CPU及其内存 device:指GPU及其内存使用nvcc,就可以编译C ...
CUDA PTX ISA阅读笔记（二）
8. 第八章指令集这一章占了整个手册的一大半(百十来页吧),主要介绍各种指令,虽然页数很多,但是大多数指令都很简单. 8.1. 指令的形式和语义描述这章就是主要描述每个PTX指令.除了指令的形式 ...
OV7670无FIFO读寄存器成功
使用stm32f103zet6 接3V3,GND SIOC 接PD0,SIOD接PD1 XCLK接PA8 SCCB.c #include "SCCB.h"#define SCCB_ ...
哔哩哔哩小甲鱼汇编语言记录一下寄存器（内存访问）两个小实验
第一个实验: 打开win10中的模拟器. 1, ,2, 2, 3,通过-t命令进行单步执行.执行完所有已经输入的程序就可以看到最后的AX BX CX寄存器内容. 最后的AX BX CX: AX=112 ...
汇编语言基础教程-寄存器
32位CPU一般包括如下寄存器 1.通用寄存器 8个32位寄存器 EAX,EBX,ECX,EDX,ESI,EDI,EBP,ESP 8个16位寄存器 AX,BX,CX,DX,SI,DI,BP,SP 8个 ...

cuda ptx 汇编语言示例：读寄存器

cuda ptx 汇编语言示例：读寄存器相关推荐

最新文章

热门文章