CUDA Dynamic Parallelism测试

news2024/9/30 23:26:15

CUDA Dynamic Parallelism测试

  • 一.知识点
  • 二.测试内容
    • 1.查看动态并行生成的PTX
    • 2.性能对比测试(测试一个向量的N次累加)
    • 3.动态并行是如何调度SM的
  • 三.查看动态并行生成的PTX
  • 四.性能对比测试
  • 五.动态并行是如何调度SM的

CUDA 动态并行（CUDA Dynamic Parallelism）是 NVIDIA 在其 CUDA 编程模型中引入的一个强大特性。它允许 GPU 上运行的内核（kernel）直接在设备端启动新的内核，而无需返回主机（CPU）进行控制。这一特性使得我们可以在 GPU 上实现更复杂、更动态的算法，提高程序的并行度和执行效率。

一.知识点

  1. CUDA 动态并行性（Dynamic Parallelism）：这是对 CUDA 编程模型的扩展，允许 CUDA 内核直接在 GPU 上创建和同步新工作。它使得在程序的任何需要的地方动态创建并行性成为可能。
  2. 减少主机和设备之间的数据传输：通过在设备上运行的线程在运行时决定启动配置，动态并行性可以减少在主机和设备之间传输执行控制和数据的需要。
  3. 数据驱动的并行工作生成：在运行时，内核可以根据数据驱动的决策或工作负载，在内核内生成依赖于数据的并行工作，动态利用 GPU 的硬件调度器和负载均衡器。
  4. 表达复杂的算法和编程模式：以前需要修改以消除递归、不规则循环结构或其他不适合单级并行性的算法，现在可以更透明地表达。
  5. 支持的计算能力：动态并行性仅支持计算能力为 3.5 及以上的设备。
  6. CUDA 执行模型的扩展：支持动态并行性的 CUDA 执行模型现在允许设备线程配置、启动新网格（grids），并在设备上对其进行隐式同步。
  7. 父子网格的关系：
    • 父线程、线程块、网格：启动新网格的实体，被称为父级。
    • 子网格：由父级启动的新网格。
    • 嵌套执行：子网格的启动和完成是正确嵌套的，父网格在其所有子网格完成之前不会被视为完成。
  8. 设备运行时的作用：提供使内核函数能够使用动态并行性的运行时系统和 API。
  9. 网格范围内的资源共享：在设备上，所有线程在网格内共享已启动的内核和 CUDA 对象。这意味着一个线程创建的流可以被网格内的任何其他线程使用。
  10. 流和事件的使用：
    • 设备上创建的流：仅在创建它们的网格范围内存在，超出该范围的行为是未定义的。
    • NULL 流的特殊性：在设备上，隐式的 NULL 流只在线程块内共享，不同线程块中的线程对 NULL 流的启动可能会并发执行。
  11. 并发执行的保证：CUDA 编程模型不保证设备上不同线程块之间的并发执行，包括父网格和子网格之间。
  12. 多 GPU 支持的限制：设备运行时不支持多 GPU；它只能在当前执行的设备上操作。
  13. 内存一致性和可见性：
    • 全局和常量内存：父子网格共享相同的全局和常量内存，但具有独立的本地和共享内存。
    • 内存一致性保证：子网格只有在启动时保证与父线程的内存视图是一致的。由于 cudaDeviceSynchronize() 的移除，父网格无法在退出前保证看到子网格的内存修改。
  14. 零拷贝内存：与全局内存具有相同的一致性保证，但内核不能在设备上分配或释放零拷贝内存。
  15. 设备运行时 API：
    • 类似于主机运行时 API：设备运行时的语法和语义与主机运行时 API 基本相同，便于代码重用。
    • 内核启动的异步性：与主机端启动相同，设备端内核启动相对于启动线程是异步的。
    • 不支持的功能：设备运行时不支持像 cudaStreamSynchronize() 和 cudaStreamQuery() 这样的 API，也不支持从设备上创建或销毁纹理和表面对象。
  16. 特殊流的使用：
    • Fire-and-Forget 流（cudaStreamFireAndForget）：用于立即调度启动，无需依赖之前的启动，无法与事件配合使用。
    • Tail Launch 流（cudaStreamTailLaunch）：用于在父网格完成后调度新的网格启动，同样无法与事件配合使用。
  17. 事件的限制：仅支持用于流间同步的 CUDA 事件，不支持 cudaEventSynchronize()、cudaEventElapsedTime() 和 cudaEventQuery() 等功能。
  18. 设备属性查询的限制：只能查询当前设备的属性，不支持在设备运行时切换设备。
  19. 全局和常量内存变量的行为：设备上的所有内核都可以读取或写入全局变量，但不能修改常量内存中的数据。
  20. 错误处理：
    • 错误代码的获取：每个线程可以通过 cudaGetLastError() 获取其生成的最后一个错误代码。
    • 错误传播：子网格中的错误（例如访问无效地址）将返回到主机。
  21. PTX 支持：CUDA 提供了底层的 PTX API，如 cudaLaunchDevice() 和 cudaGetParameterBuffer()，供需要在 PTX 级别支持动态并行性的编程语言和编译器实现者使用。
  22. 编译和链接：
    • 不需要显式包含头文件：在编译 CUDA 程序时，会自动包含设备运行时 API 的原型。
    • 设备运行时库：使用动态并行性的 CUDA 程序需要链接设备运行时静态库 libcudadevrt。
  23. 系统资源的限制和配置：
    • 启动池的大小：受限于系统资源，可以使用 cudaDeviceSetLimit() 配置启动池的大小。
    • 堆栈大小的控制：可以通过 cudaDeviceSetLimit() 设置每个 GPU 线程的堆栈大小。
  24. 内存分配注意事项：
    • 设备上的 cudaMalloc() 和 cudaFree()：在设备上调用时，与主机上的行为不同，映射到设备端的 malloc() 和 free()，受限于设备 malloc 堆大小。
    • 指针的限制：在设备上分配的内存指针不能在主机上释放，反之亦然。
  25. 线程重调度的注意事项：设备运行时可能会将线程块重新调度到不同的 SM，以更有效地管理资源，因此依赖 %smid 或 %warpid 保持不变是不安全的。
  26. ECC 错误处理：CUDA 内核中无法通知 ECC 错误，所有 ECC 错误将在整个启动树完成后在主机端报告。

二.测试内容

1.查看动态并行生成的PTX

2.性能对比测试(测试一个向量的N次累加)

  • A.调用N次Kernel
  • B.在Kernel里循环N次
  • C.使用动态并行,递归N次

3.动态并行是如何调度SM的

三.查看动态并行生成的PTX

tee dynamic_parallelism.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>

#define CHECK_CUDA(call)                      \
  do {                              \
    cudaError_t err = call;                  \
    if (err != cudaSuccess) {                 \
      std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \
      std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \
      exit(EXIT_FAILURE);                  \
    }                             \
  } while (0)


__global__ void kernel(float *iodata,int count)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  if(count>0)
  {
    iodata[tid]+=1;
    if(tid == 0)
    {
      __prof_trigger(0);
      kernel<<<gridDim.x, blockDim.x,0,cudaStreamFireAndForget >>>(iodata, count - 1);
      __prof_trigger(1);
    }
  }
}

int main(int argc,char *argv[])
{
  int deviceid=0;cudaSetDevice(deviceid); 
  int block_count=100000;
  int block_size=1024;
  int count=1000;
    
  size_t value;
  CHECK_CUDA(cudaDeviceGetLimit(&value, cudaLimitDevRuntimePendingLaunchCount));
  printf("cudaLimitDevRuntimePendingLaunchCount:%ld\n",value);

  {
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=0;
      kernel<<<block_count, block_size>>>(iodata,count);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }  
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -ptx -arch=sm_86 -rdc=true  \
    -o dynamic_parallelism.ptx dynamic_parallelism.cu \
    -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
cat dynamic_parallelism.ptx
  • 输出
.version 8.1
.target sm_86
.address_size 64

        // .globl       _Z6kernelPfi
.extern .func  (.param .b64 func_retval0) __cudaCDP2GetParameterBufferV2
(
        .param .b64 __cudaCDP2GetParameterBufferV2_param_0,
        .param .align 4 .b8 __cudaCDP2GetParameterBufferV2_param_1[12],
        .param .align 4 .b8 __cudaCDP2GetParameterBufferV2_param_2[12],
        .param .b32 __cudaCDP2GetParameterBufferV2_param_3
);
.extern .func  (.param .b32 func_retval0) __cudaCDP2LaunchDeviceV2
(
        .param .b64 __cudaCDP2LaunchDeviceV2_param_0,
        .param .b64 __cudaCDP2LaunchDeviceV2_param_1
);

.visible .entry _Z6kernelPfi(
        .param .u64 _Z6kernelPfi_param_0,
        .param .u32 _Z6kernelPfi_param_1
)
{
  .reg .pred      %p<4>;
  .reg .f32       %f<3>;
  .reg .b32       %r<11>;
  .reg .b64       %rd<8>;
  
  ld.param.u64    %rd2, [_Z6kernelPfi_param_0];
  ld.param.u32    %r2, [_Z6kernelPfi_param_1];
  mov.u32         %r1, %ntid.x;
  setp.lt.s32     %p1, %r2, 1;
  @%p1 bra        $L__BB0_5;

  mov.u32         %r3, %ctaid.x;
  mov.u32         %r4, %tid.x;
  mad.lo.s32      %r5, %r3, %r1, %r4;
  cvta.to.global.u64      %rd3, %rd2;
  mul.wide.u32    %rd4, %r5, 4;
  add.s64         %rd5, %rd3, %rd4;
  ld.global.f32   %f1, [%rd5];
  add.f32         %f2, %f1, 0f3F800000;
  st.global.f32   [%rd5], %f2;
  setp.ne.s32     %p2, %r5, 0;
  @%p2 bra        $L__BB0_5;

  // begin inline asm
  pmevent         0;
  // end inline asm
  mov.u32         %r6, %nctaid.x;
  mov.u32         %r7, 1;
  mov.u64         %rd6, _Z6kernelPfi;
  mov.u32         %r8, 0;
  { // callseq 0, 0
    .reg .b32 temp_param_reg;
    .param .b64 param0;
    st.param.b64    [param0+0], %rd6;
    .param .align 4 .b8 param1[12];
    st.param.b32    [param1+0], %r6;
    st.param.b32    [param1+4], %r7;
    st.param.b32    [param1+8], %r7;
    .param .align 4 .b8 param2[12];
    st.param.b32    [param2+0], %r1;
    st.param.b32    [param2+4], %r7;
    st.param.b32    [param2+8], %r7;
    .param .b32 param3;
    st.param.b32    [param3+0], %r8;
    .param .b64 retval0;
    call.uni (retval0), __cudaCDP2GetParameterBufferV2,(param0,param1,param2,param3);
    ld.param.b64    %rd1, [retval0+0];
  } // callseq 0
  setp.eq.s64     %p3, %rd1, 0;
  @%p3 bra        $L__BB0_4;
  add.s32         %r9, %r2, -1;
  st.u64  [%rd1], %rd2;
  st.u32  [%rd1+8], %r9;
  mov.u64         %rd7, 4;
  { // callseq 1, 0
    .reg .b32 temp_param_reg;
    .param .b64 param0;
    st.param.b64    [param0+0], %rd1;
    .param .b64 param1;
    st.param.b64    [param1+0], %rd7;
    .param .b32 retval0;
    call.uni (retval0),__cudaCDP2LaunchDeviceV2,(param0,param1); 
    ld.param.b32    %r10, [retval0+0];
  } // callseq 1
$L__BB0_4:
  // begin inline asm
  pmevent         1;
  // end inline asm
$L__BB0_5:
  ret;
}

四.性能对比测试

tee dynamic_parallelism.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>
#include <chrono>
#include <thread>

#define CHECK_CUDA(call)                      \
  do {                              \
    cudaError_t err = call;                  \
    if (err != cudaSuccess) {                 \
      std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \
      std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \
      exit(EXIT_FAILURE);                  \
    }                             \
  } while (0)

__global__ void case_0(float *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  iodata[tid]+=1;
}

__global__ void case_1(float *iodata,int count=100)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  for(int i=0;i<count;i++)
  {
      iodata[tid]+=1;
  }
}

__global__ void case_2(float *iodata,int count,bool is_block)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  if(count>0)
  {
    iodata[tid]+=1;
    if(is_block) __syncthreads(); // 同步所有线程
    if(tid == 0)
      case_2<<<gridDim.x, blockDim.x,0>>>(iodata, count - 1,is_block);
  }
}

__global__ void case_3(float *iodata,int count)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  if(count>0)
  {
    iodata[tid]+=1;
    if(tid == 0)
      case_3<<<gridDim.x, blockDim.x,0,cudaStreamFireAndForget >>>(iodata, count - 1);
  }
}

template <typename F>
void TIMEIT(F const &f,cudaStream_t &stream,cudaEvent_t &start_ev,cudaEvent_t&stop_ev)
{ 
    f(stream); 
    CHECK_CUDA(cudaDeviceSynchronize());
    auto start = std::chrono::high_resolution_clock::now();
    cudaEventRecord(start_ev, stream); 
    f(stream); 
    cudaEventRecord(stop_ev, stream); 
    CHECK_CUDA(cudaEventSynchronize(stop_ev)); 
    auto end = std::chrono::high_resolution_clock::now(); 
    std::chrono::duration<double> diff = end - start; 
    float milliseconds = 0; 
    cudaEventElapsedTime(&milliseconds, start_ev, stop_ev); 
    printf("E2E:%7.2fms Kernel:%7.2fms\n",diff.count()*1000,milliseconds);
}

int main(int argc,char *argv[])
{
  int deviceid=0;cudaSetDevice(deviceid); 
  int block_count=100000;
  int block_size=1024;
  int count=1000;
  
  cudaStream_t stream;
  cudaStreamCreate(&stream);

  cudaEvent_t start_ev, stop_ev;
  cudaEventCreate(&start_ev);
  cudaEventCreate(&stop_ev);
  
  size_t value;
  CHECK_CUDA(cudaDeviceGetLimit(&value, cudaLimitDevRuntimePendingLaunchCount));
  printf("cudaLimitDevRuntimePendingLaunchCount:%ld\n",value);

  {//Host循环Lanuch count次
      printf(" ----------------- case 0 ----------------- \n");
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=0;
      TIMEIT([&](cudaStream_t &stream)-> void {for(int i=0;i<count;i++){case_0<<<block_count, block_size,0,stream>>>(iodata);}},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }
  {//Kernel内循环加count次
      printf(" ----------------- case 1 ----------------- \n");
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=0;
      TIMEIT([&](cudaStream_t &stream)-> void {case_1<<<block_count, block_size>>>(iodata,count);},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }
  {//Kernel通过动态并行,递归count次,每次同步线程块
      printf(" ----------------- case 2 ----------------- \n");
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=0;
      TIMEIT([&](cudaStream_t &stream)-> void {case_2<<<block_count, block_size>>>(iodata,count,true);},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }
  {//Kernel通过动态并行,递归count次,每次不需要同步线程块
      printf(" ----------------- case 3 ----------------- \n");
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=0;
      TIMEIT([&](cudaStream_t &stream)-> void {case_2<<<block_count, block_size>>>(iodata,count,false);},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }  
  {//Kernel通过动态并行,递归count次,每次不需要同步线程块,cudaStreamFireAndForget
      printf(" ----------------- case 4 ----------------- \n");
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=0;
      TIMEIT([&](cudaStream_t &stream)-> void {case_3<<<block_count, block_size>>>(iodata,count);},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }  
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -rdc=true -lineinfo \
    -o dynamic_parallelism dynamic_parallelism.cu \
    -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./dynamic_parallelism
  • 输出
cudaLimitDevRuntimePendingLaunchCount:2048
 ----------------- case 0 -----------------
E2E:84655.40ms Kernel:84654.82ms   #Host循环Lanuch count次
2000.000000 2000.000000
 ----------------- case 1 -----------------
E2E:  85.26ms Kernel:  82.59ms     #Kernel内循环加count次
2000.000000 2000.000000
 ----------------- case 2 -----------------
E2E:85089.90ms Kernel:85088.64ms   #Kernel通过动态并行,递归count次,每次同步线程块
2000.000000 2000.000000
 ----------------- case 3 -----------------
E2E:84906.45ms Kernel:84904.73ms   #Kernel通过动态并行,递归count次,每次不需要同步线程块
2000.000000 2000.000000
 ----------------- case 4 -----------------
E2E:84757.05ms Kernel:84755.00ms   #Kernel通过动态并行,递归count次,每次不需要同步线程块,cudaStreamFireAndForget
2000.000000 2000.000000

五.动态并行是如何调度SM的

tee dynamic_parallelism.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>
#include <chrono>
#include <thread>

#define CHECK_CUDA(call)                      \
  do {                              \
    cudaError_t err = call;                  \
    if (err != cudaSuccess) {                 \
      std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \
      std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \
      exit(EXIT_FAILURE);                  \
    }                             \
  } while (0)

__global__ void child_kernel(unsigned int *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  unsigned int smid;
  asm volatile("mov.u32 %0, %smid;" : "=r"(smid));  
  iodata[tid]=smid;
}

__global__ void main_kernel(unsigned int *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  if(tid == 0) child_kernel<<<28,1024,0>>>(iodata);
}

int main(int argc,char *argv[])
{
  int deviceid=0;cudaSetDevice(deviceid); 
  int count=28*1024;
  
  unsigned int last_smid=-1;
  {
      unsigned int *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,count*sizeof(unsigned int),cudaHostAllocDefault));
      for(int i=0;i<count;i++) iodata[i]=0;
      main_kernel<<<1, 1>>>(iodata);
      CHECK_CUDA(cudaDeviceSynchronize());
      for(int i=0;i<count;i++) 
      {
        if(iodata[i]!=last_smid)
        {
          printf("tid:%06d smid:%04d\n",i,iodata[i]);
          last_smid=iodata[i];
        }
      }
      CHECK_CUDA(cudaFreeHost(iodata));
  }
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -rdc=true -lineinfo \
    -o dynamic_parallelism dynamic_parallelism.cu \
    -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./dynamic_parallelism
  • 输出
tid:000000 smid:0000
tid:001024 smid:0002
tid:002048 smid:0004
tid:003072 smid:0006
tid:004096 smid:0008
tid:005120 smid:0010
tid:006144 smid:0012
tid:007168 smid:0014
tid:008192 smid:0016
tid:009216 smid:0018
tid:010240 smid:0020
tid:011264 smid:0022
tid:012288 smid:0024
tid:013312 smid:0026
tid:014336 smid:0001
tid:015360 smid:0003
tid:016384 smid:0005
tid:017408 smid:0007
tid:018432 smid:0009
tid:019456 smid:0011
tid:020480 smid:0013
tid:021504 smid:0015
tid:022528 smid:0017
tid:023552 smid:0019
tid:024576 smid:0021
tid:025600 smid:0023
tid:026624 smid:0025
tid:027648 smid:0027

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.coloradmin.cn/o/2180999.html

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈，一经查实，立即删除！

相关文章

python14_运算符复合赋值

python14_运算符复合赋值

复合赋值缩写 A 7 B 3 C "hello" D "world" E True F False# 加法赋值运算符,7 3 10 def add1(a, b):a b # 等同于a a breturn a# 字符串加法赋值运算符,hello world helloworld def add2(c, d):c d # 等同于字符串拼接,c c dreturn c# …
阅读更多...
尚乐代驾重做

尚乐代驾重做

微信小程序客户端登陆 一 总体流程 前段携带code发送请求给后端&#xff0c;后端利用微信小程序相关api解析这个code&#xff0c;获取能唯一认证登陆身份的openid。接着先到数据库查询是否有这个id&#xff0c;如果没有就保存用户信息实现注册&#xff0c;并返回token到前端…
阅读更多...
百度百科 X-Bk-Token 算法还原

百度百科 X-Bk-Token 算法还原

声明 本文章中所有内容仅供学习交流，严禁用于商业用途和非法用途，否则由此产生的一切后果均与作者无关，若有侵权，请私信我立即删除！ 文章目录 声明案例地址参数分析X-Bk-Token算法追踪X-Bk-Token后缀算法还原c 值跟踪与算法还原往期逆向文章推荐最近太忙了，博客摆烂了好…
阅读更多...
算法题——层序遍历（一层按一组输出）

算法题——层序遍历（一层按一组输出）

输入[1,2,3,4,5] 输出[[1],[2,3],[4,5]] (按照一层输出为一组) 思路&#xff1a; 使用队列&#xff0c;在队列内层再加入一层for循环&#xff0c;每层的个数就是队列中当前队列的个数。 python&#xff1a; from collections import dequeclass Solution(Object):def levelO…
阅读更多...
VUE a-table 动态拖动修改列宽＋固定列

VUE a-table 动态拖动修改列宽＋固定列

实现效果 实现思路 自定义表头&#xff0c;在标题后面加两个标签&#xff0c;分别用来显示拖拽图标&#xff08;cursor: col-resize&#xff09;&#xff0c;和蓝色标记线&#xff08;有的时候鼠标移动过程中不一定会在表内&#xff0c;这个时候不显示图标&#xff0c;只显示蓝…
阅读更多...
独立样本t检验及其案例分析

独立样本t检验及其案例分析

作者简介&#xff1a;热爱数据分析&#xff0c;学习Python、Stata、SPSS等统计语言的小高同学~个人主页&#xff1a;小高要坚强的博客当前专栏&#xff1a;SPSS本文内容&#xff1a;独立样本t检验及其案例分析作者“三要”格言&#xff1a;要坚强、要努力、要学习 目录 一、引…
阅读更多...
2024版Clion debug无法查看函数内数组内容 解决办法

2024版Clion debug无法查看函数内数组内容 解决办法

参考Clion debug查看数组中的内容&#xff0c;新版本有所变化 众所周知&#xff0c;进入函数的数组debug不显示内容&#xff0c;解决办法如下&#xff1a; 在Evaluate expression中输入 *(int(*)[10])(数组名)
阅读更多...
硬件设计基础之闲聊千兆以太网

硬件设计基础之闲聊千兆以太网

一、千兆以太网的介绍 常见的以太网有几种形式&#xff0c;10Mbps、100MHz、1000Mbps、10Gbps等&#xff0c;当然还有更高的&#xff0c;接触的少些&#xff0c;暂且不聊。 10Mbps、100Mbps&#xff0c;多见于单片机使用&#xff0c;比如STM32/GD32 1000Mbps、10Gbps&#x…
阅读更多...
React学习笔记（3.0）

React学习笔记（3.0）

classnames优化类名控制 classnames是一个简单的JS库&#xff0c;可以非常方便的通过条件动态控制class类名的显示。 安装classnames&#xff1a; npm i classnames 导入&#xff1a; import classNames from classnames <div className{classNames(box3,{box2:11})}&g…
阅读更多...
ESXI识别USB设备

ESXI识别USB设备

步骤&#xff1a; 插入usb设备到服务器。关闭虚拟机&#xff0c;添加USB控制器&#xff0c;根据U盘选择usb 3.0控制器&#xff0c;再添加usb设备如果usb设备灰色&#xff0c;进入主机打开SSH。使用xshell进行连接&#xff0c;运行以下命令&#xff1a; ESXI识别USB设备 - 插入…
阅读更多...
浏览器插件的标准项目结构通常包括以下几个目录和文件

浏览器插件的标准项目结构通常包括以下几个目录和文件

浏览器插件的标准项目结构通常包括以下几个目录和文件&#xff1a; my-extension/ ├── manifest.json // 插件的元数据和配置 ├── background.js // 背景脚本 ├── content_scripts/ // 内容脚本目录 │ └── content.js // 内容脚本…
阅读更多...
如何使用ssm实现基于在线开放课程的Web前端设计与实现+vue

如何使用ssm实现基于在线开放课程的Web前端设计与实现+vue

TOC ssm746基于在线开放课程的Web前端设计与实现vue 绪论 1.1 选题背景 当人们发现随着生产规模的不断扩大&#xff0c;人为计算方面才是一个巨大的短板&#xff0c;所以发明了各种计算设备&#xff0c;从结绳记事&#xff0c;到算筹&#xff0c;以及算盘&#xff0c;到如今…
阅读更多...
基于ssm的宠物领养管理系统的设计与实现 （含源码+sql+视频导入教程）

基于ssm的宠物领养管理系统的设计与实现 （含源码+sql+视频导入教程）

&#x1f449;文末查看项目功能视频演示获取源码sql脚本视频导入教程视频 1 、功能描述 基于ssm的宠物领养管理系统3拥有两种角色 管理员&#xff1a;用户管理、管理员管理、宠物管理、领养管理、评论管理、团队活动管理、志愿者申请管理、领养列表 用户&#xff1a;查看各种…
阅读更多...
51单片机应用开发（进阶）---数码管+按键+蜂鸣器（电磁炉显示模拟）

51单片机应用开发（进阶）---数码管+按键+蜂鸣器（电磁炉显示模拟）

实现目标 1、加强数码管、按键的学习&#xff0c;实现数码显示变量数据&#xff08;四位数的显示&#xff09;&#xff1b; 2、4位数码2个按键无源蜂鸣器实现模拟电磁炉功率调节及显示&#xff1b; 一、内容描述 功能描述&#xff1a;1、开机显示电磁炉功率300&#xff0c;每…
阅读更多...
Percona Monitoring and Management

Percona Monitoring and Management

Percona Monitoring and Management (PMM)是一款开源的专用于管理和监控MySQL、MongoDB、PostgreSQL
阅读更多...
Java入门3——操作符+String

Java入门3——操作符+String

在入门2中忘了提 String 的事情了&#xff0c;所以这篇就放在开头啦&#xff0c;很有用 话不多说&#xff0c;开始正题~ 一、String 引用数据类型之——String 1.字符串的拼接 在Java中&#xff0c;如果要把两句话合并到一句话的时候&#xff0c;其实是很简单的&#xff0c;…
阅读更多...
0109 图解多线程死锁问题

0109 图解多线程死锁问题

死锁场景 &#x1f92a;举例场景&#xff1a;两个憨憨Tom和Sam去西餐厅吃牛排&#xff0c;桌子上只有一把刀和一把叉&#xff0c;Tom先拿到了叉子&#xff0c;Sam拿到了刀。只有同时拿到刀叉才能吃牛排&#xff0c;于是两个憨憨陷入如下的僵局。 这个场景中&#xff0c;就存在…
阅读更多...
信号处理快速傅里叶变换(FFT)的学习

信号处理快速傅里叶变换(FFT)的学习

FFT是离散傅立叶变换的快速算法&#xff0c;可以将一个信号变换到频域。有些信号在时域上是很难看出什么特征的&#xff0c;但是如果变换到频域之后&#xff0c;就很容易看出特征了。这就是很多信号分析采用FFT变换的原因。另外&#xff0c;FFT可以将一个信号的频谱提取出来&am…
阅读更多...
可以白嫖PPT模板的6个网站，赶紧收藏

可以白嫖PPT模板的6个网站，赶紧收藏

推荐6个PPT模板网站&#xff0c;免费下载&#xff0c;绝对的高质量&#xff0c;赶紧收藏&#xff01; 1、菜鸟图库 ppt模板免费下载|ppt背景图片 - 菜鸟图库 菜鸟图库网有非常丰富的免费素材&#xff0c;像设计类、办公类、自媒体类等素材都很丰富。PPT模板种类很多&#xff0…
阅读更多...
高中教辅汇总【35GB】

高中教辅汇总【35GB】

文章目录 一、资源概览二、资源亮点三、获取方式 一、资源概览 这份教辅资源汇总&#xff0c;精心搜集了高中各学科的海量教辅资料&#xff0c;总容量高达35GB&#xff0c;覆盖了语文、数学、英语、物理、化学、生物、历史、地理、政治等所有必修及选修科目。从基础知识点到难…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023