ipadmini2降级10.3.3是啥版本

Posted 2023-04-23

rA9发布了一款全新A7降级工具LeetDown

支持A7处理器完美降级ios 10.3.3

支持的型号5s、iPad Air1、iPad mini2，可以从任何iOS版本降级到10.3.3

使用该工具并不需要越狱才能使用，只是需要MacOS系统，玄烨品果推荐Mac 10.15.7，白苹果、黑苹果

2021.2.24日更新日志：

添加 A6 SoC 支持以降级到 iOS 8.4.1。
为任何版本的 macOS 11+ 添加官方 M1 支持。
极大地提高了应用程序和漏洞利用的可靠性。
修复了 iPad 上的引导链上传错误。

注意：Apple Silicon 支持是实验性的，在某些条件下可能不起作用，请参阅自述文件以获取详细说明

下载地址：

提取码：356655

iPad mini 2完美降级10.3.3教程（iPad Air/5S通用）
1，下载上面的工具，最新2.2版本。

2，5S、iPad mini2、iPad Air进入DFU模式。

3，打开工具LeetDown，点击Select 10.3.3 iPSW，选择已经下载好的10.3.3的固件。

4，点击Browse for an IPSW 选择已经下载的10.3.3固件，或者选择Download an iPSW在线下载固件

5，接下来点击Downgrade，将会降级10.3.3版本，清空设备所有数据，点击Ok

6，提示是否保存Blob，点击Skip跳过。

7，此时工具开始降级，等待自动降级。

8，设备会出现如下画面，开始走进度条，表示已经开始降级了，耐心等待。

9，降级完成后，重新设置iPad mini2 ，降级完成。

注意：第4步调用漏洞后，不是每次都能成功，具体看个人的电脑情况，失败十几次也是有可能。如果手机屏幕变绿色就是成功了。不成功的反应是会自动开机，那么就需要重新进入DFU模式再试。

如果打开软件提示文件损坏的解决办法：

•打开 /应用程序/ 实用工具/终端

•将LeetDown 拖动到/应用程序/

复制这条命令: xattr -cr /Applications/Leetdown.app，并粘贴到终端中。

•按Enter

•完成！
文章标签： 5S降级 IPAD AIR1降级 MINI2降级 降级

为您推荐

HFZ Activator Ramdisk 工具绕过物主与锁定，支持16.3

TiggerRamDisk绕过激活界面，支持最新iOS16.2绕过

LU RAMDISK AIO 绕过iOS15激活软件，支持14.0~ 15.7.1版本

iRemovalPro完美绕iOS 16.3三网信号，支持重启iCloud

minaRActivator完美绕iOS 16.2信号，支持打电话4G流量iCloud登录

DkRamdisk 绕过iOS9.3~10.3.3,支持完美激活重启iTunes同步
发表评论
要发表评论，您必须先登录。 参考技术A 答案如下：ipadmini2降级10.3.3是亚太大陆版本，与此同时，车辆完成了较为完善的底盘包围，前后保险杠位置我们都能够看到不同功能的金属材料来完成防刮蹭，而侧边虽然不是运动样式，仍旧完成了独立侧裙包围，在对于车身的整体防护与通过性改善方面是有益的。

PTX 中的变量“已降级”是啥意思？

【中文标题】PTX 中的变量“已降级”是啥意思？

【英文标题】：What does it mean when a variable "has been demoted" in the PTX?PTX 中的变量“已降级”是什么意思？

【发布时间】：2021-10-03 04:13:11

【问题描述】：

在我的 CUDA 内核的函数体中，我有几个固定大小的 __shared__ 数组变量。当我查看其中一个数组的已编译 PTX 代码 (SM 7.5) 时，我看到一条评论说：

// my_kernel(t1 p1, t2 p2)::my_variable has been demoted

... 此行出现在 PTX 中的 .global 行中，就在编译内核本身之前。然后，在内核中，我得到：

// demoted variable
.shared .align 4 .b8 my_kernel(t1 p1, t2 p2)::my_variable [1234];

我的问题：

这样的变量是通过什么方式“降级”的？不是按照我的要求定义的吗？ 在什么情况下这些变量会被“降级”？

注意事项：

我使用的是 CUDA 11.2。 我从我的 PTX 中引用了去错线。实际名称是_ZZ8blahblah...。 我看到这个“降级”的数组变量要么是二维固定大小的数组，要么是具有结构的元素类型（例如struct unsigned short data[2]; ）；也许这有某种关系。

【问题讨论】：

@talonmies: 1. 我假设这具有标准含义；但根据你的建议，我会生成一个 MRE。 2. “全局->共享->全局”循环是什么意思？我从一个共享变量开始。

@RobertCrovella：这些变量可能无法优化到寄存器中（它们太大并且在编译时不知道访问索引）；但让我看看那个 MRE。

【参考方案1】：

根据here 的讨论，这似乎是基于共享变量范围是否可以限制为单个函数（即单个内核）。即使是我看过的具有共享使用的非常复杂的内核函数也会降级共享变量。

这是一个降级和不降级的简单示例。

未降级：

$ vi t1.cu
$ cat t1.cu
__shared__ float s[32];
__global__ void k(float *  my_ptr)
  s[threadIdx.x] = threadIdx.x;
  *my_ptr = s[threadIdx.x];

__global__ void k1(float *  my_ptr)
  s[threadIdx.x] = 0.0f;
  *my_ptr = s[threadIdx.x];

$ nvcc -ptx t1.cu
$ cat t1.ptx
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-29373293
// Cuda compilation tools, release 11.2, V11.2.67
// Based on NVVM 7.0.1
//

.version 7.2
.target sm_52
.address_size 64

        // .globl       _Z1kPf
.shared .align 4 .b8 s[128];

.visible .entry _Z1kPf(
        .param .u64 _Z1kPf_param_0
)

        .reg .f32       %f<2>;
        .reg .b32       %r<5>;
        .reg .b64       %rd<3>;


        ld.param.u64    %rd1, [_Z1kPf_param_0];
        cvta.to.global.u64      %rd2, %rd1;
        mov.u32         %r1, %tid.x;
        cvt.rn.f32.u32  %f1, %r1;
        shl.b32         %r2, %r1, 2;
        mov.u32         %r3, s;
        add.s32         %r4, %r3, %r2;
        st.shared.f32   [%r4], %f1;
        st.global.f32   [%rd2], %f1;
        ret;


        // .globl       _Z2k1Pf
.visible .entry _Z2k1Pf(
        .param .u64 _Z2k1Pf_param_0
)

        .reg .b32       %r<6>;
        .reg .b64       %rd<3>;


        ld.param.u64    %rd1, [_Z2k1Pf_param_0];
        cvta.to.global.u64      %rd2, %rd1;
        mov.u32         %r1, %tid.x;
        shl.b32         %r2, %r1, 2;
        mov.u32         %r3, s;
        add.s32         %r4, %r3, %r2;
        mov.u32         %r5, 0;
        st.shared.u32   [%r4], %r5;
        st.global.u32   [%rd2], %r5;
        ret;

降级：

$ cat t1.cu
__global__ void k(float *  my_ptr)
  __shared__ float s[32];
  s[threadIdx.x] = threadIdx.x;
  *my_ptr = s[threadIdx.x];

__global__ void k1(float *  my_ptr)
  __shared__ float s[32];
  s[threadIdx.x] = 0.0f;
  *my_ptr = s[threadIdx.x];

$ nvcc -ptx t1.cu
$ cat t1.ptx
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-29373293
// Cuda compilation tools, release 11.2, V11.2.67
// Based on NVVM 7.0.1
//

.version 7.2
.target sm_52
.address_size 64

        // .globl       _Z1kPf
// _ZZ1kPfE1s has been demoted
// _ZZ2k1PfE1s has been demoted

.visible .entry _Z1kPf(
        .param .u64 _Z1kPf_param_0
)

        .reg .f32       %f<2>;
        .reg .b32       %r<5>;
        .reg .b64       %rd<3>;
        // demoted variable
        .shared .align 4 .b8 _ZZ1kPfE1s[128];

        ld.param.u64    %rd1, [_Z1kPf_param_0];
        cvta.to.global.u64      %rd2, %rd1;
        mov.u32         %r1, %tid.x;
        cvt.rn.f32.u32  %f1, %r1;
        shl.b32         %r2, %r1, 2;
        mov.u32         %r3, _ZZ1kPfE1s;
        add.s32         %r4, %r3, %r2;
        st.shared.f32   [%r4], %f1;
        st.global.f32   [%rd2], %f1;
        ret;


        // .globl       _Z2k1Pf
.visible .entry _Z2k1Pf(
        .param .u64 _Z2k1Pf_param_0
)

        .reg .b32       %r<6>;
        .reg .b64       %rd<3>;
        // demoted variable
        .shared .align 4 .b8 _ZZ2k1PfE1s[128];

        ld.param.u64    %rd1, [_Z2k1Pf_param_0];
        cvta.to.global.u64      %rd2, %rd1;
        mov.u32         %r1, %tid.x;
        shl.b32         %r2, %r1, 2;
        mov.u32         %r3, _ZZ2k1PfE1s;
        add.s32         %r4, %r3, %r2;
        mov.u32         %r5, 0;
        st.shared.u32   [%r4], %r5;
        st.global.u32   [%rd2], %r5;
        ret;

顺便说一句，确实存在 PTX 生成器可以完全删除共享变量的情况，但这与这里的问题没有直接关系。

【讨论】：

嗯，在我的例子中，这些变量是先验限制在单个函数中的——因为它们是在内核函数的主体中定义的。所以我的情况本质上是第二个k1。但是那么......那里发生的“降级”是什么？函数范围的共享变量出现在 PTX 和 CUDA 函数中。

降级是将变量的范围从隐式全局范围转换为该函数的局部范围。在这种情况下，该变量显然已经在本地范围内，但前端处理中也可能发生促销活动。您可能希望阅读我链接的整个主题。