程序员经验分享-hwhale

reduceSum记录

2023-05-16

reduceSum 算子实现版本记录总结

第一版

__global__ void reduceSum(float *input, float *result, const int size) {
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  if (x >= size) {
    return;
  }

  int len = blockDim.x / 2;
  while (len > 0) {
    if (threadIdx.x < len) {
      input[x] += input[x + len];
    }

    len /= 2;
    __syncthreads();
  }

  if (threadIdx.x == 0) {
    result[blockIdx.x] = input[blockIdx.x];
  }
}

主要思想是在一个block内每次规约启动左边一半的线程，每一个线程完成当前规模的一次原地规约，随后规模减半，直到为一，此时
该块的规约结果就放在线程0管理的地址处。

这个kernel没有线程束分化的问题，因为每次激活一半的线程，同一个线程束内的线程要么同时启动，要么不参与运算，不存在一部分
线程选择A分支，另一部分线程选择B分支的情况。

第二版

__global__ void reduceSumSharedMemory(float *input, float *result,
                                      const int size) {
  extern __shared__ float tempArray[];
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  if (x >= size) {
    return;
  }

  // prepare data
  tempArray[threadIdx.x] = input[x];
  __syncthreads();

  int len = blockDim.x / 2;
  while (len > 0) {
    if (threadIdx.x < len) {
      tempArray[threadIdx.x] += tempArray[threadIdx.x + len];
    }

    len /= 2;
    __syncthreads();
  }

  // write internal result back to global memory
  if (threadIdx.x == 0) {
    result[blockIdx.x] = tempArray[threadIdx.x];
  }
  __syncthreads();
}

主要思想是由于全局内存的存取效率相对低于对共享内存的存取，所以第一步可以先将数据从全局内存读取到共享内存中，随后在共享内存中
进行规约。

主要结果如下：

dataseize1031024*1024
Durationusecond744.58
SOL DRAM%26.22
SOL L1/TEX Cache%94.54
SOL L2 Cache%9.81
SM [%]%75.04

相比于第一版并没有时间上的进步，但是提高了L1 Cache的利用率，主要原因在于算子的计算密度并不高，主要限制在于内存带宽的利用率。虽然
在共享内存中规约肯定会比从全局内存中规约要快，但是算子把数据从全局内存读取到共享内存这一步效率并不高，即使能够很快的完成共享内存
规约，也会因为等待数据时间过长导致没有任何进步。

第三版

__global__ void reduceSumTwoLoads(float *input, float *result, const int size) {
  extern __shared__ float tempArray[];

  int x = blockIdx.x * blockDim.x * 2 + threadIdx.x;
  if (x >= size) {
    return;
  }

  // prepare data
  tempArray[threadIdx.x] = input[x] + input[x + blockDim.x];
  __syncthreads();

  int len = blockDim.x / 2;
  while (len > 0) {
    if (threadIdx.x < len) {
      tempArray[threadIdx.x] += tempArray[threadIdx.x + len];
    }

    len /= 2;
    __syncthreads();
  }

  // write internal result back to global memory
  if (threadIdx.x == 0) {
    result[blockIdx.x] = tempArray[threadIdx.x];
  }
}

这个kernel的主要改进点在于，kernel执行的时候每次规约都用一半的线程不活跃，造成线程资源的浪费，因此，在上一步的基础上每次读取两个数据，并进行规约
结果如下：

dataseize1031024*1024
Durationusecond372.22
SOL DRAM%52.08
SOL L1/TEX Cache%91.3
SOL L2 Cache%19.09
SM [%]%73.88

相比于上一个kernel，这次的改进导致时间接近减半，内存利用率翻倍。最初改进的理由是避免线程的浪费，从SM利用率上看的确实现了这一目标，但是kernel的性能
改进主要得益于内存带宽的利用率提高。同一个线程每次读取的数据增加了，全局层面上会有更多的访存被合并，提高了数据读取到共享内存的效率，减少了数据等待的时间
自然提高了性能。

第四版

__device__ void wrapReduce(volatile float *input, int id) {
  input[id] += input[id + 32];
  input[id] += input[id + 16];
  input[id] += input[id + 8];
  input[id] += input[id + 4];
  input[id] += input[id + 2];
  input[id] += input[id + 1];
}
__global__ void reduceSumUnrollLastWrap(float *input, float *result, int size) {
  extern __shared__ float tempArray[];

  int x = blockIdx.x * blockDim.x * 2 + threadIdx.x;
  if (x >= size) {
    return;
  }

  // prepare data
  tempArray[threadIdx.x] = input[x] + input[x + blockDim.x];
  __syncthreads();

  int len = blockDim.x / 2;
  while (len > 32) {
    if (threadIdx.x < len) {
      tempArray[threadIdx.x] += tempArray[threadIdx.x + len];
    }

    len /= 2;
    __syncthreads();
  }

  // last wrap
  if (threadIdx.x < 32) {
    wrapReduce(tempArray, threadIdx.x);
  }

  // write internal result back to global memory
  if (threadIdx.x == 0) {
    result[blockIdx.x] = tempArray[threadIdx.x];
  }
}

这一版的改进出发点在于，当规约规模小于32时只需要一个线程束就能实现所有的功能，而GPU上线程的基本调度单位是线程束，在同一个线程束内
的所有线程使用SIMD的方式同步执行。因此，当规模小于等于32时对循环进行展开，直接在一个束内进行求解，这样的好处是避免了所有线程在规模
小于等于32时执行无效的for循环和if判断指令，提高了线程的利用效率。

结果如下：

dataseize1031024*1024
Durationusecond247.52
SOL DRAM%78.40
SOL L1/TEX Cache%59.03
SOL L2 Cache%28.79
SM [%]%52.44

从结果看，随着内存带宽利用率的提高，算子耗时进一步减少，主要原因在于通过展开最后一个线程束，减少了线程的无效指令执行时间，从而间接
增加了线程访存的时间。

第五版

template <int threadNumber>
__device__ void reduceOnSharedMempry(float *tempArray, int id) {
  // manually unroll max to 1024 threads

  if (threadNumber >= 1024) {
    if (threadIdx.x < 512) {
      tempArray[threadIdx.x] += tempArray[threadIdx.x + 512];
    }
    __syncthreads();
  }

  if (threadNumber >= 512) {
    if (threadIdx.x < 256) {
      tempArray[threadIdx.x] += tempArray[threadIdx.x + 256];
    }
    __syncthreads();
  }

  if (threadNumber >= 256) {
    if (threadIdx.x < 128) {
      tempArray[threadIdx.x] += tempArray[threadIdx.x + 128];
    }
    __syncthreads();
  }

  if (threadNumber >= 128) {
    if (threadIdx.x < 64) {
      tempArray[threadIdx.x] += tempArray[threadIdx.x + 64];
    }
    __syncthreads();
  }

  // last wrap
  if (threadIdx.x < 32) wrapReduce<128>(tempArray, threadIdx.x);
}

template <int threadNumber>
__global__ void reduceSumUnrollIterations(float *input, float *result) {
  __shared__ float tempArray[threadNumber];
  int x = blockIdx.x * blockDim.x * 2 + threadIdx.x;

  // prepare data
  tempArray[threadIdx.x] = input[x] + input[x + blockDim.x];
  __syncthreads();

  reduceOnSharedMempry<threadNumber>(tempArray, threadIdx.x);

  // write internal result back to global memory
  if (threadIdx.x == 0) {
    result[blockIdx.x] = tempArray[threadIdx.x];
  }
}

这一版的改进思路延续上一版的循环展开，既然最后一个束展开后可以减少for循环和if判断指令的消耗，那么在明确线程块大小的前提下也可以通过
循环展开减少循环指令和if指令的消耗。

结果如下：

dataseize1031024*1024
Durationusecond239.42
SOL DRAM%80.80
SOL L1/TEX Cache%60.81
SOL L2 Cache%29.63
SM [%]%54.01

和上一版只展开最后一个束相比，推理时间进一步减少，这种循环展开带来的性能提升主要是减少了不必要的指令执行，增加了访存指令的密度，间接
提高了访存时间和内存带宽利用率。

第六版

template <int threadNumber>
__global__ void reduceSumUnrollAndLoads4(float *input, float *result) {
  __shared__ float tempArray[threadNumber];
  int x = blockIdx.x * blockDim.x * 4 + threadIdx.x;

  // loads four elements into register
  float temp = input[x] + input[x + blockDim.x];
  float temp2 = input[x + blockDim.x * 2] + input[x + blockDim.x * 3];
  // write results into shared memory
  tempArray[threadIdx.x] = temp + temp2;
  __syncthreads();

  // reduce on shared memory
  reduceOnSharedMempry<threadNumber>(tempArray, threadIdx.x);

  // write internal result back to global memory
  if (threadIdx.x == 0) {
    result[blockIdx.x] = tempArray[threadIdx.x];
  }
}

这个kernel的优化思路延续第三版，既然同一个线程读取两个数据会提高内存带宽的利用率，那么增加更多的数据读取应该能进一步增带宽利用率，
这一版是一个线程块处理连续的四个数据块。

结果如下：

dataseize3210241024
DRAM Frequencycycle/nsecond6.78
SM Frequencycycle/nsecond1.36
Elapsed Cyclescycle319030
Memory [%]%88.28
SOL DRAM%88.28
Durationusecond234.05
SOL L1/TEX Cache%32.74
SOL L2 Cache%31.75
SM Active Cyclescycle315,965.44
SM [%]%29.06

这一版的内存利用率达到了88.28，推理时间相比于之前的版本也有了更多的进步，提升的主要原因就在于进一步地增加了可供调度的访存指令的数量，
增加了访存指令的密度，从而提高内存带宽的利用率。

第七版 

template <int threadNumber>
__global__ void reduceSumUnrollAndLoads8(float *input, float *result) {
  __shared__ float tempArray[threadNumber];
  int x = blockIdx.x * blockDim.x * 8 + threadIdx.x;

  // loads four elements into register
  float temp = input[x] + input[x + blockDim.x];
  float temp2 = input[x + blockDim.x * 2] + input[x + blockDim.x * 3];
  float temp3 = input[x + blockDim.x * 4] + input[x + blockDim.x * 5];
  float temp4 = input[x + blockDim.x * 6] + input[x + blockDim.x * 7];
  // write results into shared memory
  tempArray[threadIdx.x] = temp + temp2 + temp3 + temp4;
  __syncthreads();

  // reduce on shared memory
  reduceOnSharedMempry<threadNumber>(tempArray, threadIdx.x);

  // write internal result back to global memory
  if (threadIdx.x == 0) {
    result[blockIdx.x] = tempArray[threadIdx.x];
  }
}

这个kernel的优化思路延续第六版，既然增加独立的访存指令能够提高内存带宽利用率，那么就可以继续增加，这个版本一个块需要处理原来8倍的block。

结果如下：

dataseize3210241024
DRAM Frequencycycle/nsecond6.79
SM Frequencycycle/nsecond1.36
Elapsed Cyclescycle317,563
Memory [%]%88.62
SOL DRAM%88.62
Durationusecond232.74
SOL L1/TEX Cache%21.31
SOL L2 Cache%47.55
SM Active Cyclescycle314,594.09
SM [%]%19.46

一次读取8个数据并进行规约进一步的提高了内存带宽的利用率，从而减少了推理时间，由于一个线程需要处理原来八倍的数据量，所需的block数量大大降低，
所以SM的利用率反而降低了

第八版

template <int threadNumber>
__global__ void reduceSumUnrollAndLoads16LinearAccessMemory(float *input,
                                                            float *result) {
  __shared__ float tempArray[threadNumber];
  int x = blockIdx.x * blockDim.x * 16 + threadIdx.x * 16;

  // loads 16 elements into register
  float temp = input[x] + input[x + 1];
  float temp2 = input[x + 2] + input[x + 3];
  float temp3 = input[x + 4] + input[x + 5];
  float temp4 = input[x + 6] + input[x + 7];
  float temp5 = input[x + 8] + input[x + 9];
  float temp6 = input[x + 10] + input[x + 11];
  float temp7 = input[x + 12] + input[x + 13];
  float temp8 = input[x + 14] + input[x + 15];

  // write results into shared memory
  tempArray[threadIdx.x] =
      temp + temp2 + temp3 + temp4 + temp5 + temp6 + temp7 + temp8;
  __syncthreads();

  // reduce on shared memory
  reduceOnSharedMempry<threadNumber>(tempArray, threadIdx.x);

  // write internal result back to global memory
  if (threadIdx.x == 0) {
    result[blockIdx.x] = tempArray[threadIdx.x];
  }
}

这一版尝试继续增加单个线程处理的数据量。

结果如下：

dataseize3210241024
DRAM Frequencycycle/nsecond6.80
SM Frequencycycle/nsecond1.37
Elapsed Cyclescycle523,348
Memory [%]%53.75
SOL DRAM%53.75
Durationusecond382.98
SOL L1/TEX Cache%97.92
SOL L2 Cache%19.16
SM Active Cyclescycle518,682.57
SM [%]%8.85

从结果看并没有达到理想的效果，一方面是因为这种通过增加单个线程负载的方式来提高内存带宽利用率已经接近极限了，在上一个版本中的
内存利用率提升已经很小了再增加负载效果有限。

性能衰退的原因可能是因为寄存器资源不足，导致单个线程不能按照设定的逻辑将数据存放
到寄存器中然后直接规约。实际上这一系列优化的主要思路就是通过给单个线程分配更多的寄存器资源和增加更多的内存访问实现对内存带宽
的利用率增加，这一方法的一个限制就在于SM上的寄存器资源，当资源不足时，寄存器就无法放置这么多全局内存数据，转而将数据存放在全局内存
导致性能下降。

第九版 

template <int threadNumber>
__global__ void reduceSumUnrollAndLoads8LinearAccessMemory(float *input,
                                                           float *result) {
  __shared__ float tempArray[threadNumber];
  int x = blockIdx.x * blockDim.x * 8 + threadIdx.x * 8;

  // loads eight elements into register
  float temp = input[x] + input[x + 1];
  float temp2 = input[x + 2] + input[x + 3];
  float temp3 = input[x + 4] + input[x + 5];
  float temp4 = input[x + 6] + input[x + 7];
  // write results into shared memory
  tempArray[threadIdx.x] = temp + temp2 + temp3 + temp4;
  __syncthreads();

  // reduce on shared memory
  reduceOnSharedMempry<threadNumber>(tempArray, threadIdx.x);

  // write internal result back to global memory
  if (threadIdx.x == 0) {
    result[blockIdx.x] = tempArray[threadIdx.x];
  }
}

这一版在第七版的基础上进行优化，主要考虑是第七版读取八个全局内存数据的时候间隔很大，这样可能会造成分散的内存访问，降低内存带宽利用率。

结果如下：

dataseize3210241024
DRAM Frequencycycle/nsecond6.79
SM Frequencycycle/nsecond1.36
Elapsed Cyclescycle318,004
Memory [%]%88.48
SOL DRAM%88.48
Durationusecond232.12
SOL L1/TEX Cache%89.35
SOL L2 Cache%40.24
SM Active Cyclescycle314,906.72
SM [%]%19.14

和第七版相比，推理时间有一点进步，但是并不是和预想的那样的进步，主要原因在于，块内的合并内存访问极大的提高了L1 Cache的利用率，所以减少了
规约操作的时间，但是块内的合并内存访问并没有提高全局的内存合并访问，从结果上看甚至降低了内存利用率，全靠L1 Cache的提高平衡了一些性能。

第十版

template <int threadNumber>
__device__ void wrapReduceByShuffle(volatile float *input, int id) {
  if (threadNumber >= 64) input[id] += input[id + 32];

  // read data into thread register
  int sum = input[id];

  // Use XOR mode to perform butterfly reduction
  sum += __shfl_xor_sync(0xffffffff, sum, 16, 32);
  sum += __shfl_xor_sync(0xffffffff, sum, 8, 32);
  sum += __shfl_xor_sync(0xffffffff, sum, 4, 32);
  sum += __shfl_xor_sync(0xffffffff, sum, 2, 32);
  sum += __shfl_xor_sync(0xffffffff, sum, 1, 32);

  // write register value to shared memory
  if (id == 0) *input = sum;
}

这一版优化的出发点是，在展开最后一个束的时候，第四版是将数据从共享内存读入寄存器，规约完成后再写回全局内存，这就存在重复的数据读写和运算指令，
通过一次读数据后直接使用shuffle指令即可实现线程束内部的规约，减少重复的数据读写和运算指令消耗。

结果差不多，几乎没有区别。

分析原因，访问寄存器只需要一个周期，访问共享内存是1~20个周期，理论上寄存器应该更快，但是算子的主要瓶颈仍然在于全局内存访问效率，即使规约操作
变快了还是需要等待数据，因此这样的优化意义不大。另一方面，由于对共享内存的访问不存在bank conflict 的问题，本身访问共享内存的带宽就足够大，和
访问寄存器差异不大。

第十一版

__global__ void reduceSumByWrap(float *input, float *result) {
  __shared__ float tempArray[32];

  int wrapNum = blockDim.x / 32;
  int wrapId = threadIdx.x >> 5;
  int laneId = threadIdx.x & 31;
  volatile float *basePtr = input + wrapId * 32;

  // read data into registers and perform wrap reduce
  float value = *(basePtr + laneId);
  value += __shfl_xor_sync(0xffffffff, value, 16, 32);
  value += __shfl_xor_sync(0xffffffff, value, 8, 32);
  value += __shfl_xor_sync(0xffffffff, value, 4, 32);
  value += __shfl_xor_sync(0xffffffff, value, 2, 32);
  value += __shfl_xor_sync(0xffffffff, value, 1, 32);
  if (laneId == 0) {
    tempArray[wrapId] = value;
  }
  __syncthreads();

  // reduce on shared memory
  if (wrapId == 0) {
    if (laneId < wrapNum)
      value = tempArray[laneId];
    else
      value = 0;

    value += __shfl_xor_sync(0xffffffff, value, 16, 32);
    value += __shfl_xor_sync(0xffffffff, value, 8, 32);
    value += __shfl_xor_sync(0xffffffff, value, 4, 32);
    value += __shfl_xor_sync(0xffffffff, value, 2, 32);
    value += __shfl_xor_sync(0xffffffff, value, 1, 32);
  }
  if (laneId == 0) result[blockIdx.x] = value;
}

这一版的优化思路建立在对之前的kernel的数据流动分析基础上，以目前最好的第九版为例：数据首先从全局内存流向寄存器，然后从寄存器流向共享内存，这就是
第一步的数据导入；在接下来的基于共享内存的规约中，每进行一次规约活跃线程数量减半，造成线程资源的浪费。

为了解决这两个问题，可以使用基于束的规约方式：第一步中，每个线程束获取到全局内存数据后直接使用shuffle指令进行束内规约，然后将结果写入共享内存，这样即可
解决第一个问题；在基于共享内存的规约过程中，每个线程束处理32个数据，读入数据后进行束内规约，然后将结果写入共享内存。使用这样的规约方式既能够优化数据
流动，又可以优化线程资源使用。

结果如下：

dataseize3210241024
DRAM Frequencycycle/nsecond6.79
SM Frequencycycle/nsecond1.36
Elapsed Cyclescycle2,203,537
Memory [%]%14.70
SOL DRAM%0.00
Durationusecond1600
SOL L1/TEX Cache%15.76
SOL L2 Cache%5.09
SM Active Cyclescycle2,201,373.46
SM [%]%14.70

结果看来，这个kernel性能非常差，并没有达到预期效果。主要原因在于，当线程束读取全局内存数据的时候采用的是束内连续读取，这样的方式在微观上能够一次读取32个
数据，但是从全局上并不能导致合并的内存读取，因此内存的利用率十分低下。

第十二版

__global__ void reduceSumByWrapV2(float *input, float *result, size_t size) {
  __shared__ float tempArray[32];

  int wrapNum = blockDim.x / 32;
  int wrapId = threadIdx.x >> 5;
  int laneId = threadIdx.x & 31;

  float value = 0;
  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < size;
       i += blockDim.x * gridDim.x) {
    value += input[i];
  }

  // read data into registers and perform wrap reduce
  value += __shfl_xor_sync(0xffffffff, value, 16, 32);
  value += __shfl_xor_sync(0xffffffff, value, 8, 32);
  value += __shfl_xor_sync(0xffffffff, value, 4, 32);
  value += __shfl_xor_sync(0xffffffff, value, 2, 32);
  value += __shfl_xor_sync(0xffffffff, value, 1, 32);
  if (laneId == 0) {
    tempArray[wrapId] = value;
  }
  __syncthreads();

  // reduce on shared memory
  if (wrapId == 0) {
    if (laneId < wrapNum)
      value = tempArray[laneId];
    else
      value = 0;

    value += __shfl_xor_sync(0xffffffff, value, 16, 32);
    value += __shfl_xor_sync(0xffffffff, value, 8, 32);
    value += __shfl_xor_sync(0xffffffff, value, 4, 32);
    value += __shfl_xor_sync(0xffffffff, value, 2, 32);
    value += __shfl_xor_sync(0xffffffff, value, 1, 32);
  }
  if (laneId == 0) result[blockIdx.x] = value;
}

这版本参考了Nvidia 的博客，在上一个版本的基础上，使用grid-stride的方式读取全局内存而不是基于束的全局内存读取，其余操作不变。

结果如下：

dataseize3210241024
DRAM Frequencycycle/nsecond6.79
SM Frequencycycle/nsecond1.36
Elapsed Cyclescycle318,146
Memory [%]%88.46
SOL DRAM%88.46
Durationusecond232.12
SOL L1/TEX Cache%59.30
SOL L2 Cache%40.09
SM Active Cyclescycle315,140.16
SM [%]%19.43

从结果上看，使用了grid-stride方式读取全局内存后，改进的wrap-based kernel性能能够达到之前通过增加内存事务一样的效果。这个kernel的好处在于，
使用grid-stride方式读取全局内存使得kernel具有可扩展性，之前的方式需要指定thread的数量刚好匹配的处理的数据的数量，如果要处理多个数据就需要
改变kernel调用的方式。使用新的wrap-based, grid-stride读取方式的kernel能够使用相同的代码调用，不需要复杂的调用函数。

还没有想明白的问题

完全循环展开为什么128效果最好，256，512效果不好，

如何设定block和grid

block中线程的数量应该大于一个SM中最大线程数/最大block数，并且选择SM中最大线程数的约数。由于block是在SM上运行，grid中block的数量最好是SM个数
的倍数。

其它收获

在优化reduce算子的过程中纠正了一些错误的认知，比如”使用共享内存一定会更快“，”一个block内的线程访问连续的地址会提高全局内存访问效率“，更重要的
是在分析算子的时候能够从数据流动的方式分析算子的性能，通过分析数据的流动分析算子是否有冗余的数据搬运或者逻辑分支。

本文内容由网友自发贡献，版权归原作者所有，本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容，请联系:hwhale#tublm.com(使用前将#替换为@)

reduceSum

reduceSum记录 的相关文章

  • 边缘计算的解决方案大集合

    自今年2月的巴塞罗那世界移动通信大会召开以来 xff0c 边缘计算无疑是C位出道 xff0c 爆发释放在人们的视野中 xff0c 成为今年业界最热门的领域之一 顺着5G的东风 xff0c 边缘计算的诞生成为历史必然 xff0c 整个行业都在

  • 计网(笔记版)---外部网络路由协议之BGP协议

  • js中?. 、?? 、??=的用法及含义

    1 可选链运算符 是不是经常遇到这样的错误 TypeError Cannot read properties of null reading 39 xxx 39 引入可选链就是为了解决这个问题 const person 61 id 1 na

  • LXC与Docker介绍

    文章目录 LXCLUX是什么LXC常用命令LXC的使用 Docker容器虚拟化和传统虚拟化的区别Linux NamespacesCGroupsdoeker基本概念docker容器编排 LXC LUX是什么 LXC xff08 LinuX C

  • 树莓派使用CLASH的代理安装软件

    为什么要使用代理 github系列域名不能访问 xff0c curl一键安装用不了 打开CLASH允许局域网功能 树莓派终端登陆 方法一 xff1a 1 编辑 etc profile文件 sudo nano etc profile 2 在最

  • Linux下解决高并发socket最大连接数限制，tcp默认1024个连接

    linux获取TCP连接数 方法一 xff1a admin 64 zabbix ss ant awk 39 NR gt 1 a 1 43 43 END for b in a print b a b 39 ESTAB 535 TIME WAI

  • vncserver的详细配置

    原文地址 xff1a vncserver的详细配置 作者 xff1a OpenTech 1 首先要配置的是服务端 A 确认服务器端是否安装了vncserver 使用rpm qa vnc命令如果收到如下信息说明已经安装了vncserver x

  • Rman备份中常见的问题

    1 xff1a ORA 01031insufficient privileges gpasswd d oracle dba 将oracle移除出dba组 查看oracle属性 uid 61 500 oracle gid 61 500 oin

  • 弱监督学习-snorkel

    1 什么是弱监督学习 弱监督问题旨在研究通过较弱的监督信号来构建预测模型 xff0c 即在少量的标注样本上学习建模 xff0c 达到大量样本上同样的效果 弱监督学习主要分为三类 不确切监督 xff08 inexact supervision

  • 如何用3000元搞定一年100M点对点专线

    温馨提示 xff1a 阅读本文需要先阅读或温习格物资讯早先发布过的 奇葩物花生壳出品蒲公英VPN组网路由 2015年11月格物资讯发布了花生壳打洞路由器蒲公英的试用报告 xff0c 提到在鹏博士接入前提下 xff0c 做P2P组网实现快速的

  • PNETLAB中可以导入的交换机、防火墙等设备镜像

    在网上找了很久 xff0c 想要找到一个设备镜像的下载 xff0c 发现网上全都是一些对于PNET本体安装的炒冷饭 不过经过一个下午的寻找 xff0c 最终在B站一个UP 64 real半吊子工程师 22年的视频里找到了相关的下载平台连接

  • 在Keras中，TimeDistributed层的作用是什么？

    在Keras中 xff0c TimeDistributed层的作用是什么 xff1f 关键词 xff1a python xff0c machine learning xff0c keras xff0c neural network xff0

  • 理解1D、2D、3D卷积神经网络的概念

    目录 引言二维CNN Conv2D一维CNN Conv1D三维CNN Conv3D总结 引言 当我们说卷积神经网络 xff08 CNN xff09 时 xff0c 通常是指用于图像分类的二维CNN 但是 xff0c 现实世界中还使用了其他两

  • 解决vncserver看不到桌面的问题

    解决vncserver看不到桌面的问题 主要参考这里 xff1a http zhidao baidu com link url 61 7Btj0KsV5b986dydoOpElKDpSwriaruP4jxWY6f6pG3Ota kcQbdV

  • 深入理解 keras 中 Dense 层参数

    目录 引言深入理解 Dense 层的用法查看参数输入尺寸输出尺寸示例 xff1a 用法完整示例示例一 最小网络示例二 xff1a 多维度数据示例三 xff1a 特殊情况 xff0c 待讨论 附录 引言 大家或许已经对深度学习不陌生了 不管是

  • 如何在Keras中使用数据生成器（data generators）的详细示例

    目录 动机讲解以前的情况小提示数据产生器Keras脚本 可运行实例结论 动机 您是否曾经不得不加载一个非常消耗内存的数据集 xff0c 以至于希望魔术能够无缝地解决这一问题 xff1f 大型数据集正日益成为我们生活的一部分 xff0c 因为

  • 绘制 x+y+z=1 图像

    简单手动推导 matlab 绘制 figure fimplicit3 64 x y z abs x 43 abs y 43 abs z 4 figure fimplicit3 64 x y z x 43 y 43 z 1 ref How d

  • yolov5 test.py val.py detec.py 区别在哪里呢?

    yolov5 test py val py detec py 区别在哪里呢 用户在训练数据的时候必须使用 train py 来进行 数据训练和验证 xff0c 但我很难理解 detect py 和 test py 之间的区别 应该在一个数据

随机推荐

  • window、wsl2

    主要涉及到的点包括 xff1a window 11 下 wsl2 的安装 xff1b 在WSL上使用NVIDIA SDK manager给Jetson烧录系统 xff1b 加载 USB 设备 xff1b ref xff1a https de

  • Jetson基础知识

    Jetson基础知识 在网上安装软件 看到到处都在讲 Jetson 的版本 看到很多名词也不是很懂 特别是那些缩写 本文就是根据查找到的资料进行统一的解释说明 L4T L4T 代表 Linux for Tegra 这是一个专门为NVIDIA

  • Ubuntu桌面

    Ubuntu桌面 1 桌面系统 安装桌面的时候 xff0c Ubuntu18系统默认用的时候 gnome 桌面 但是有两种类型 xff1a ubuntu desktop 和 ubuntu gnome desktop ubuntu deskt

  • 常用各类数据集

    原文链接 xff1a http homepages inf ed ac uk rbf CVonline Imagedbase htm CVonline xff1a 图像数据库 Google直译的结果 xff0c 希望对大家有帮助 按主题索引

  • vnc无画面的解决方法

    网上看到了很多vnc配置文件 xff0c 例如安装xfce后 xff0c 只需要几行代码 xff0c 就可以正常显示画面 而自己实际操作的时候 xff0c 发现怎么弄都是白屏 xff0c 没有窗口 xff0c 非常困惑 经过了一段时间的折腾

  • 打造基于 VNC 的 Ubuntu 20.04 的远程桌面

    VNC server B站视频 打造基于 VNC 的Ubuntu 20 04远程桌面 https www bilibili com video BV1nV41147dt spm id from 61 333 999 0 0 1 安装 VNC

  • “Top-down”---至顶向下的设计方法

    Top down 至顶向下的设计方法 曾经看到有人说 xff0c 人活着的过程就是在不断地解决问题的过程 我觉得这句话很有道理 xff0c 从年幼时的牙牙学语 xff0c 到学习阶段的各种作业 xff0c 当然还有各种编程难题 xff0c

  • linux与其他操作系统文件共享方法

    大家好 xff0c 我是加摩斯 xff0c 觉得文章有帮助的小伙伴 xff0c 记得一键三连哟 xff5e 申明 xff1a 原创 xff0c 转载前请与我沟通 前言 xff1a 我将Linux文件共享的方法分为以下几类 xff1a 1 x

  • 树莓派4B安装 Ubuntu20.04 + ROS Noetic 踩坑记录

    写在开头 本文记录了给一台树莓派4B配置ROS开发环境的全部操作过程 xff0c 和在实际操作过程中遇到的各种问题的解决方法 xff0c 希望对有相同需求的小伙伴们有帮助 本文的目标是直接在树莓派上开发ROS xff0c 因而在安装时与目标

  • python输出格式—format方法

    format方法 format 功能很强大 xff0c 它把字符串当成一个模板 xff0c 通过传入的参数进行格式化 xff0c 并且使用大括号 作为特殊字符代替C语言中的 位置映射 print 34 34 format 39 123 39

  • python 字符串详解（附案例）

    目录 什么是python字符串 如何表示一个字符串 字符串的访问 print xff08 xff09 索引 for循环 打印字符 end 切片split 范围选择符 n m 字符串长度 使用len xff08 xff09 函数 replac

  • SQL语言多表查询

    多表查询的基础是单表查询 xff0c 一些基本的语句都在单表查询中介绍过了 目录 主外键相等链接 xff08 两个表 xff09 join链接 xff08 两个表 xff09 子查询 xff08 两个表 xff09 自链接 xff08 两个

  • resize 详细讲解 C++

    resize 函数 是一个替换字符串长度的函数 xff0c 有两个重载函数 xff0c 第一个参数都是替换之后的大小 xff0c 第二个为替换的字母 include lt iostream gt include lt string gt u

  • C++内联函数 如何使用、优缺点

    将inline写在函数之前就成了内联函数 内联函数适用于 频繁调用的小函数 xff08 语句简单 xff09 内联函数的优点 内联函数是为了取代C语言中的宏而存在的 因为宏比较不好写 xff08 被括号和分号支配的恐惧 xff09 宏不能调

  • C++ 类和对象（上） 访问限定符 域操作符 this指针 类的定义

    访问限定符 public xff1a 在类外可以直接被访问 private xff1a 在类外不可以直接被访问 protected xff1a 在类外不可以直接被访问 class 的默认访问权限为private xff0c strict 的

  • C++ 类与对象（中）构造函数 析构函数

    构造函数 什么是构造函数 C 43 43 是基于C的基础上的 xff0c 但是为了提升C 语言 xff0c C 43 43 在类中引用了构造函数 构造函数是为了将类初始化 构造函数的特点 1 名字与类名相同 2 创建类类型对象时由编译器自动

  • mysql: failed to connect to localhost:3306

    mysql failed to connect to localhost 3306 这个问题是因为mysql的特定服务没有开启 xff0c 客户端通过网络发起请求自然是登陆不上的 解决方法 打开windows服务管理器 xff0c 找到my

  • 命令行模式安装VisualStudio

    命令行模式安装VisualStudio 今天远程安装VisualStudio2019遇到神奇bug xff0c installer运行后完全透明 xff0c 根本无法界面上配置安装 xff0c 还好找到了命令行模式 首先下载 exe 这里选

  • Ubuntu 设置中文语言环境

    点击右上角的齿轮形状的按钮 xff0c 然后选择System Settings 点击进入 会显示有语言需要安装 点击install 再次当前用户输入密码 xff0c 显示安装过程中 点击 install remove languages 选

  • reduceSum记录

    reduceSum 算子实现版本记录总结 第一版 global void reduceSum float input float result const int size int x 61 blockIdx x blockDim x 43

热门标签