train_gpt2_fp32.cu - layernorm_forward_kernel3

源码

__global__ void layernorm_forward_kernel3(float* __restrict__ out, float* __restrict__ mean, float* __restrict__ rstd,const float*  __restrict__ inp, const float*  __restrict__ weight,const float* __restrict__ bias, int N, int C) {cg::thread_block block = cg::this_thread_block();cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);int idx = blockIdx.x * warp.meta_group_size() + warp.meta_group_rank();if(idx >= N) {return;}// the row of input that this group of threads is responsible forconst float* x = inp + idx * C;// meanfloat sum = 0.0f;for (int i = warp.thread_rank(); i < C; i += warp.size()) {sum += x[i];}sum = cg::reduce(warp, sum, cg::plus<float>{});float m = sum / C;if(warp.thread_rank() == 0 && mean != nullptr) {__stcs(mean + idx, m);}// rstdsum = 0.0f;for (int i = warp.thread_rank(); i < C; i += warp.size()) {float diff = x[i] - m;sum += diff * diff;}sum = cg::reduce(warp, sum, cg::plus<float>{});float s = rsqrtf(sum / C + 1e-5f);if(warp.thread_rank() == 0 && rstd != nullptr) {__stcs(rstd + idx, s);}// final normalization and scaling by weight/biasfloat* o = out + idx * C;for (int c = warp.thread_rank(); c < C; c += warp.size()) {// load and store using the .cs "streaming" hint to the compiler,// indicating that this data will not be reused soon, and can be streamed through the caches// this allows the threads to get more cache-hits for the (shared) weight and bias parametersfloat n = s * (__ldcs(x+c) - m);__stcs(o+c, n * weight[c] + bias[c]);}
}

注释

/*** @brief 层归一化前向传播内核函数** 执行层归一化的前向传播操作，并输出归一化后的结果。** @param out 归一化后的输出数组* @param mean 均值数组* @param rstd 归一化后的标准差数组* @param inp 输入数组* @param weight 权重数组* @param bias 偏置数组* @param N 样本数量* @param C 通道数量*/

/*** 在给定输入数据上执行层归一化操作的GPU内核函数。* * @param out 输出数据的指针。* @param mean 计算出的平均值的指针，如果为nullptr，则不计算和保存平均值。* @param rstd 计算出的标准差的倒数的指针，如果为nullptr，则不计算和保存标准差的倒数。* @param inp 输入数据的指针。* @param weight 权重参数的指针，用于对输出数据的每个元素进行缩放。* @param bias 偏置参数的指针，用于对输出数据的每个元素增加偏置。* @param N 输入数据的批处理大小。* @param C 输入数据的每个样本的特征数量。* * 该函数首先计算输入数据的每个特征维度的平均值，然后计算该维度的标准差的倒数，* 最后对输入数据进行归一化、缩放和偏置添加，生成输出数据。*/
__global__ void layernorm_forward_kernel3(float* __restrict__ out, float* __restrict__ mean, float* __restrict__ rstd,const float*  __restrict__ inp, const float*  __restrict__ weight,const float* __restrict__ bias, int N, int C) {cg::thread_block block = cg::this_thread_block();cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);int idx = blockIdx.x * warp.meta_group_size() + warp.meta_group_rank();if(idx >= N) {return;}// 计算负责处理的输入行const float* x = inp + idx * C;// 计算平均值float sum = 0.0f;for (int i = warp.thread_rank(); i < C; i += warp.size()) {sum += x[i];}sum = cg::reduce(warp, sum, cg::plus<float>{});float m = sum / C;if(warp.thread_rank() == 0 && mean != nullptr) {__stcs(mean + idx, m);}// 计算标准差的倒数sum = 0.0f;for (int i = warp.thread_rank(); i < C; i += warp.size()) {float diff = x[i] - m;sum += diff * diff;}sum = cg::reduce(warp, sum, cg::plus<float>{});float s = rsqrtf(sum / C + 1e-5f);if(warp.thread_rank() == 0 && rstd != nullptr) {__stcs(rstd + idx, s);}// 执行最终的归一化、缩放和偏置添加float* o = out + idx * C;for (int c = warp.thread_rank(); c < C; c += warp.size()) {// 使用.cs提示加载和存储数据，指示该数据不久将不会被重用，可以被缓存流式传输，以提高权重和偏置参数的缓存命中率float n = s * (__ldcs(x+c) - m);__stcs(o+c, n * weight[c] + bias[c]);}
}

解释

该函数是一个用于层归一化操作的CUDA内核函数，用于计算输入数据的归一化、标准化和加权求和。具体来说，它接受6个输入参数和1个输出参数，其中输入参数包括待处理的输入数据inp、缩放因子weight和偏置bias，以及待计算的均值mean和标准差rstd，输出参数为处理后的输出数据out。该函数以块的形式对输入数据进行处理，每个块处理N个数据行，每行包含C个数据列。具体步骤如下：

  1. 根据当前块的索引blockIdx.x和每个块中线程的元组数量warp.meta_group_size()计算当前线程组处理的数据行索引idx。
  2. 若idx超出输入数据行数的范围，则直接返回。
  3. 计算当前线程组处理的输入数据行的起始地址x。
  4. 分别计算该行数据的均值m和调整后标准差s，其中均值的计算采用线程块内同步的方式进行累加求和，并通过cg::reduce()函数进行归约；标准差的计算同样采用线程块内同步的方式进行累加求和，并对结果进行开方运算，并添加了一个小的常数以防止除以零。
  5. 若mean和rstd非空，则将计算得到的均值和标准差分别写入相应的输出地址。
  6. 最后，根据计算得到的均值、标准差、权重和偏置，对输入数据进行归一化、标准化和加权求和操作，并将结果写入输出地址。

该函数使用了CUDA C++的协同计算（Cooperative Groups）库cg来进行线程块级别的同步和归约操作，并使用了CUDA的流式加载和存储指令__ldcs和__stcs来优化数据加载和存储的效率。

这段 C++ 代码是一个 CUDA kernel，用于执行层归一化（Layer Normalization）的前向传播操作。层归一化是一种常用于神经网络中的技术，它用于规范化神经元的激活值，从而加速训练过程。

以下是代码的详细解释：

  1. 函数声明:
__global__ void layernorm_forward_kernel3(float* __restrict__ out, float* __restrict__ mean, float* __restrict__ rstd,const float*  __restrict__ inp, const float*  __restrict__ weight,const float* __restrict__ bias, int N, int C)
* `__global__`：这是一个 CUDA 关键字，表示这是一个 GPU 上运行的 kernel 函数。
* 函数接受以下参数：+ `out`：输出张量，存储归一化后的结果。+ `mean`：均值张量，如果提供了（即不是 `nullptr`），则存储每个输入元素的均值。+ `rstd`：规范化标准差张量，如果提供了（即不是 `nullptr`），则存储每个输入元素的规范化标准差。+ `inp`：输入张量。+ `weight`：权重张量，用于缩放归一化后的结果。+ `bias`：偏置张量，用于添加到归一化后的结果上。+ `N` 和 `C`：分别表示输入张量的高度和宽度（例如，在批量大小为 `N`、特征数为 `C` 的 2D 输入张量中）。

2. 线程索引和同步:

cg::thread_block block = cg::this_thread_block();
cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);
int idx = blockIdx.x * warp.meta_group_size() + warp.meta_group_rank();
if(idx >= N) {return;
}
* 这部分代码用于获取当前线程块的索引和线程索引，以及同步线程块中的线程。

3. 计算均值 (mean):

const float* x = inp + idx * C;
float sum = 0.0f;
for (int i = warp.thread_rank(); i < C; i += warp.size()) {sum += x[i];
}
sum = cg::reduce(warp, sum, cg::plus<float>{});
float m = sum / C;
if(warp.thread_rank() == 0 && mean != nullptr) {__stcs(mean + idx, m);
}
* 这部分代码计算输入张量中每个元素的均值。
* `cg::reduce` 用于同步线程块中的所有线程，并将它们的部分和合并为一个总和。
* 如果提供了 `mean` 张量，则将计算出的均值存储在其中。

4. 计算规范化标准差 (rstd):

sum = 0.0f;
for (int i = warp.thread_rank(); i < C; i += warp.size()) {float diff = x[i] - m;sum += diff * diff;
}
sum = cg::reduce(warp, sum, cg::plus<float>{});
float s = rsqrtf(sum / C + 1e-5f);
if(warp.thread_rank() == 0 && rstd != nullptr) {__stcs(rstd + idx, s);
}
* 这部分代码计算输入张量中每个元素的规范化标准差。
* `rsqrtf` 用于计算平方根的倒数（即规范化标准差）。
* 如果提供了 `rstd` 张量，则将计算出的规范化标准差存储在其中。

5. 最终归一化和缩放:

float* o = out + idx * C;
for (int c = warp.thread_rank(); c < C; c += warp.size()) {float n = s * (__ldcs(x+c) - m);__stcs(o+c, n * weight[c] + bias[c]);
}
* 这部分代码使用前面计算出的均值和规范化标准差对输入张量进行归一化，并使用权重和偏置进行缩放。
* `__ldcs` 和 `__stcs` 是流加载和存储操作，用于指示编译器这些数据不会很快被重用，可以通过缓存流进行传输，从而提高缓存命中率。

在 CUDA 编程中，__stcs 和类似的流操作（如 __ldcs）通常用于优化内存访问模式，特别是在处理全局内存时。这些操作提示编译器使用特定的内存访问模式，以避免缓存冲突并提高性能。然而，它们并不是必需的，并且在某些情况下可能不会被使用。

在提供的代码示例中，__stcs 用于将结果写入全局内存。如果您不使用 __stcs，而是使用普通的赋值操作（例如 *o = value;），CUDA 运行时仍然能够正确地将结果写入全局内存。但是，这可能会错过一些性能优化机会，因为编译器不会自动应用与流操作相关的优化。

是否使用 __stcs 取决于您的具体需求和性能目标。如果您对性能有严格要求，并且您的代码在全局内存访问方面遇到了瓶颈，那么使用流操作可能有助于提升性能。另一方面，如果您的代码已经足够快，或者流操作的使用过于复杂，那么您可能选择不使用它们。

请注意，__stcs 和其他类似的流操作是特定于某些 CUDA 编程模型或库的，并不是 CUDA 的标准 API 的一部分。因此，它们的使用可能取决于您所使用的具体编程环境和工具链。

总的来说，__stcs 可以在某些情况下用于优化性能，但并不是必需的。您可以根据自己的需求和目标来决定是否使用它。如果您对性能没有特殊要求，或者不熟悉这些流操作的使用，那么您可以选择不使用它们。

在 AMD GPU 编程中，没有直接对应于 NVIDIA CUDA 中的 __stcs 这样的流操作指令。AMD GPU 编程主要使用 OpenCL、HIP（Heterogeneous-compute Interface for Portability）或者 AMD 的专有工具如 AMD GPUOpen。

在 AMD GPU 上优化内存访问和进行数据传输，您通常会依赖不同的机制和工具：

  1. OpenCL 内存对象: 在 OpenCL 中，您可以使用不同的内存对象（如缓冲区对象）来管理 GPU 上的内存，并通过异步数据传输（如 clEnqueueWriteBuffer 和 clEnqueueReadBuffer）来优化数据传输。

  2. HIP: 如果您在使用 CUDA 代码并且希望将其移植到 AMD GPU，HIP 可以帮助您实现这一点。HIP 提供了一个类似于 CUDA 的 API，它会自动将 CUDA 调用映射到相应的 AMD GPU 调用。在 HIP 中，您可能会使用类似于 CUDA 的内存管理函数，但不会有 __stcs 这样的特定指令。

  3. AMD GPUOpen: AMD 提供了一套开源库和工具，用于在 AMD GPU 上进行高效编程。这些库可能提供了特定的内存管理和优化策略，但这些策略可能不同于 CUDA 的流操作。

  4. 内存访问模式: 无论使用哪种编程模型，优化内存访问模式都是关键。这包括确保数据对齐、减少内存碎片、使用适当的内存布局以及避免不必要的内存传输。

  5. 同步和调度: 在 GPU 编程中，同步和调度是重要的考虑因素。AMD GPU 编程模型通常提供了一些机制来同步不同内核之间的执行和内存访问，这对于实现高效的内存操作至关重要。

  6. 性能分析工具: 使用 AMD GPU 的性能分析工具可以帮助您识别潜在的瓶颈和优化机会，这些工具可以提供关于内存访问模式、数据传输和内核执行的有用信息。

综上所述，尽管 AMD GPU 编程没有直接的 __stcs 替代指令，但您仍然可以通过使用适当的内存管理策略、优化内存访问模式以及利用提供的编程模型和工具来在 AMD GPU 上实现高效的内存操作。

单元测试

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <float.h>
#include <string.h>#include <cuda_runtime.h>
#include <cublas_v2.h>
#include <cublaslt_v2.h>#include "utils.h"
#include "testing_utils.h"
#include "gpt2_encoder_decoder.h"#define MAX_T 20
#define MAX_B 8extern "C" void layernorm_forward_kernel3(float* out, float* mean, float* rstd,const float* inp, const float* weight,const float* bias, int B, int T, int C);void test_layernorm_forward_kernel3() {int B = 1;int T = 2;int C = 3;size_t inp_size = B * T * C;size_t out_size = inp_size;float* d_inp;float* d_out;float* d_weight;float* d_bias;float* d_mean;float* d_rstd;float* h_inp;float* h_out;float* h_weight;float* h_bias;float* h_mean;float* h_rstd;h_inp = (float*)malloc(inp_size * sizeof(float));h_out = (float*)malloc(out_size * sizeof(float));h_weight = (float*)malloc(C * sizeof(float));h_bias = (float*)malloc(C * sizeof(float));h_mean = (float*)malloc(B * C * sizeof(float));h_rstd = (float*)malloc(B * C * sizeof(float));cudaMalloc(&d_inp, inp_size * sizeof(float));cudaMalloc(&d_out, out_size * sizeof(float));cudaMalloc(&d_weight, C * sizeof(float));cudaMalloc(&d_bias, C * sizeof(float));cudaMalloc(&d_mean, B * C * sizeof(float));cudaMalloc(&d_rstd, B * C * sizeof(float));for (int i = 0; i < inp_size; i++) {h_inp[i] = i * 1.0f;}for (int i = 0; i < C; i++) {h_weight[i] = 1.0f;h_bias[i] = 0.0f;}cudaMemcpy(d_inp, h_inp, inp_size * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_weight, h_weight, C * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_bias, h_bias, C * sizeof(float), cudaMemcpyHostToDevice);layernorm_forward_kernel3<<<B, C>>>(d_out, d_mean, d_rstd, d_inp, d_weight, d_bias, B, T, C);cudaDeviceSynchronize();cudaMemcpy(h_out, d_out, out_size * sizeof(float), cudaMemcpyDeviceToHost);cudaMemcpy(h_mean, d_mean, B * C * sizeof(float), cudaMemcpyDeviceToHost);cudaMemcpy(h_rstd, d_rstd, B * C * sizeof(float), cudaMemcpyDeviceToHost);for (int b = 0; b < B; b++) {for (int c = 0; c < C; c++) {float expected = (c * T * B + b * T + 0.5f) * 1.0f;float got = h_out[b * C + c];float tolerance = 1e-5f;if (fabs(expected - got) > tolerance) {printf("layernorm_forward_kernel3 test failed at index %d, %d\n", b, c);printf("expected %f but got %f\n", expected, got);exit(1);}}}free(h_inp);free(h_out);free(h_weight);free(h_bias);free(h_mean);free(h_rstd);cudaFree(d_inp);cudaFree(d_out);cudaFree(d_weight);cudaFree(d_bias);cudaFree(d_mean);cudaFree(d_rstd);
}int main(int argc, char *argv[]) {test_layernorm_forward_kernel3();printf("layernorm_forward_kernel3 test ok\n");return 0;
}

下面是使用Google Test框架生成的单元测试代码，用于测试layernorm_forward_kernel3函数的功能：

#include <gtest/gtest.h>
#include <cuda_runtime.h>// Define the test fixture class
class LayerNormTest : public ::testing::Test {
protected:void SetUp() override {// Allocate device memory for input, output, mean, rstd, weight, and biascudaMalloc(&inp, N * C * sizeof(float));cudaMalloc(&out, N * C * sizeof(float));cudaMalloc(&mean, N * sizeof(float));cudaMalloc(&rstd, N * sizeof(float));cudaMalloc(&weight, C * sizeof(float));cudaMalloc(&bias, C * sizeof(float));// Initialize input, weight, and bias data on the hostfloat* host_inp = new float[N * C];float* host_weight = new float[C];float* host_bias = new float[C];for (int i = 0; i < N * C; i++) {host_inp[i] = i;}for (int i = 0; i < C; i++) {host_weight[i] = 1.0f;host_bias[i] = 0.0f;}// Copy host data to devicecudaMemcpy(inp, host_inp, N * C * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(weight, host_weight, C * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(bias, host_bias, C * sizeof(float), cudaMemcpyHostToDevice);delete[] host_inp;delete[] host_weight;delete[] host_bias;}void TearDown() override {// Free device memorycudaFree(inp);cudaFree(out);cudaFree(mean);cudaFree(rstd);cudaFree(weight);cudaFree(bias);}// Test dataint N = 8;int C = 16;float* inp;float* out;float* mean;float* rstd;float* weight;float* bias;
};// Define a test case
TEST_F(LayerNormTest, Forward) {// Launch the kerneldim3 block(32, 1);dim3 grid((N + block.x - 1) / block.x, 1);layernorm_forward_kernel3<<<grid, block>>>(out, mean, rstd, inp, weight, bias, N, C);// Copy output data from device to hostfloat* host_out = new float[N * C];cudaMemcpy(host_out, out, N * C * sizeof(float), cudaMemcpyDeviceToHost);// TODO: Add assertions to verify the correctness of the outputdelete[] host_out;
}// Run the tests
int main(int argc, char** argv) {::testing::InitGoogleTest(&argc, argv);return RUN_ALL_TESTS();
}

以上是使用Google Test框架生成的单元测试代码，用于测试layernorm_forward_kernel3函数的功能。在测试过程中，需要先为函数的输入和输出分配设备内存，并在测试用例中调用函数来执行计算。然后，将输出数据从设备复制到主机，并在主机端进行断言以验证输出的正确性。

请注意，在上面的代码中，需要根据实际情况添加断言来验证输出的正确性。

优化建议

这份代码是一个用于层归一化（Layer Normalization）的CUDA内核，旨在对深度学习模型中的特征图进行标准化。它计算给定输入的均值、标准差，并应用权重和偏置进行缩放和偏移。下面我将从潜在问题和优化方向两个方面进行详细评审：

潜在问题

  1. 内存访问模式：代码在访问输入inp、输出out、权重weight和偏置bias时，假设了这些数组是连续且对齐的。对于非连续或未对齐的数据，可能会导致性能下降。建议添加对数据对齐的假设说明或检查。

  2. 数据类型限制：当前实现使用float类型进行计算，对于需要更高精度或更低精度的场景可能不适用。考虑提供类型参数化的实现。

  3. 异常处理：代码中未涉及对输入参数的有效性检查（例如N和C的值是否合法，weight和bias是否为空等）。在实际应用中，应检查这些参数以避免潜在的运行时错误。

  4. 硬编码常量：如1e-5f在计算标准差时用作偏移量，这种硬编码可能会影响算法的泛化能力或适应不同应用场景的能力。建议将其作为参数传递给内核。

优化方向

  1. 代码可读性：虽然内核实现了其功能，但代码中缺乏注释，特别是对于一些不明显的操作，如__stcs和__ldcs指令。为了提高代码的可维护性，应在关键操作处添加注释，解释其目的和预期的效果。

  2. 性能优化：尽管使用了CUDA的 Cooperative Groups 和 Thread Block Tiles 功能来改善缓存效率和降低同步开销，但还可以进一步探索以下优化点：

    • 减少内存访问：计算均值和标准差时，每个线程重复访问了输入数据x。可以通过将计算均值和方差的循环合并来减少对输入数据的访问次数。

    • 利用共享内存：考虑将权重和偏置加载到共享内存中，特别是当C较大时，这可以减少全局内存访问的开销。

    • 优化归一化计算：对rsqrtf的结果再次乘以C，可能会引入额外的精度损失。考虑是否有更稳定的算法来执行这种归一化和缩放操作。

  3. 通用性：目前的实现专注于浮点数据的处理。考虑扩展此内核以支持其他数据类型（如半精度浮点类型float16），以满足对计算效率和内存使用有更高要求的应用场景。

  4. 异步执行和流控：在实际应用中，可能需要同时执行多个CUDA内核或在不同内核之间进行数据传输。考虑使用CUDA流来异步执行此内核和其他操作，以提高整体执行效率。

以下是相应的代码修订与优化。请注意，由于涉及的修改较为复杂，一些细节可能需要进一步调整。

根据您的要求，我将针对上述提出的建议逐一进行代码优化。请注意，某些建议（如检查输入参数的有效性、对硬编码常量进行参数化）在内核代码中可能不直接体现，因为它们通常在调用内核之前于主机代码中处理。不过，我会在说明中提及这些优化点。现在让我们关注代码优化：

// 定义一个宏或者常量用于调整算法的敏感度
#define Epsilon 1e-5f__global__ void layernorm_forward_kernel3(float* __restrict__ out, float* __restrict__ mean, float* __restrict__ rstd,const float*  __restrict__ inp, const float*  __restrict__ weight,const float* __restrict__ bias, int N, int C) {cg::thread_block block = cg::this_thread_block();cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);int idx = blockIdx.x * warp.meta_group_size() + warp.meta_group_rank();// 确保索引在有效范围内if(idx >= N) {return;}const float* x = inp + idx * C;// 合并计算均值和方差的循环以减少对输入的重复访问float sum = 0.0f;float square_sum = 0.0f;for (int i = warp.thread_rank(); i < C; i += warp.size()) {float diff = x[i] - sum * (1.0f / C);sum += diff;square_sum += diff * diff;}sum = cg::reduce(warp, sum, cg::plus<float>{});square_sum = cg::reduce(warp, square_sum, cg::plus<float>{});float m = sum / C;if(warp.thread_rank() == 0 && mean != nullptr) {__stcs(mean + idx, m);}float s = rsqrtf(square_sum / C + Epsilon);if(warp.thread_rank() == 0 && rstd != nullptr) {__stcs(rstd + idx, s);}float* o = out + idx * C;for (int c = warp.thread_rank(); c < C; c += warp.size()) {float n = s * (__ldcs(x+c) - m);__stcs(o+c, n * weight[c] + bias[c]);}
}

解释和应用的优化：

  1. 内存访问模式和数据类型限制：对于内存访问模式和数据类型限制的问题，建议在调用此内核的主机代码中进行检查和处理，例如确保输入输出缓冲区是适当对齐的，以及提供支持不同类型（如double或half）的版本。

  2. 异常处理：异常处理主要涉及输入参数的有效性检查，应在调用内核之前的主机代码中完成。例如，确保N和C非负，weight和bias非空等。

  3. 硬编码常量：通过定义Epsilon宏来替换硬编码的1e-5f，使得代码更易于修改和维护。

  4. 代码可读性：在关键操作处添加了注释，以提高代码的可读性和可维护性。

  5. 性能优化：

    • 合并了计算均值和方差的循环，以减少对输入数据的重复访问。
    • 考虑到内核的特性，使用共享内存或异步执行的优化可能需要对整体架构进行较大改动，可能会超出简单优化的范畴。这些优化应根据具体应用场景的需要进行仔细设计和测试。

请注意，对于任何CUDA内核，实际性能优化往往需要针对特定硬件和具体应用场景进行细致的调整和测试。以上优化建议提供了一个出发点，但可能需要根据实际情况进行进一步的优化。此优化后的代码版本仅供参考。

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

如若内容造成侵权/违法违规/事实不符，请联系长河编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

iLogtail 社区开源之夏活动来了！

iLogtail 社区开源之夏活动来了！

作者&#xff1a;玄飏 在这个充满活力的夏日&#xff0c;随着阳光一同灿烂的是开源精神的光辉与创新的火花。iLogtail 社区高兴地宣布&#xff0c;我们正式加入开源之夏 2024 的行列&#xff0c;诚邀每一位怀揣梦想与激情的学生开发者&#xff0c;共同开启一场探索技术前沿、贡…
阅读更多...
Coursera吴恩达深度学习专项课程01: Neural Networks and Deep Learning 学习笔记 Week 01

Coursera吴恩达深度学习专项课程01: Neural Networks and Deep Learning 学习笔记 Week 01

Week 01 of Neural Networks and Deep Learning Course Certificate 本文是学习 https://www.coursera.org/learn/neural-networks-deep-learning 这门课的笔记 Course Intro 文章目录 Week 01 of Neural Networks and Deep Learning[0] Welcome to the Deep Learning Spec…
阅读更多...
Apollo9.0 Control模块算法源码学习

Apollo9.0 Control模块算法源码学习

阅读更多...
Spring Boot | Spring Boot 整合 “RabbitMQ“ ( 消息中间件 ) 实现

Spring Boot | Spring Boot 整合 “RabbitMQ“ ( 消息中间件 ) 实现

目录: Spring Boot 整合 "RabbitMQ" ( 消息中间件 )实现 &#xff1a;一、Spring Boot 整合 整合实现 : Publish/Subscribe ( 发布订阅 ) 工作模式 ( "3种"整合实现方式 )1.1 基于"API"的方式 ( 实现 Publish/Subscribe "发布订阅"工作…
阅读更多...
redis抖动问题导致延迟或者断开的处理方案

redis抖动问题导致延迟或者断开的处理方案

目录&#xff1a; 1、使用背景2、redis重试机制3、redis重连机制4、其他一些解决redis抖动问题方案 1、使用背景 客户反馈文件偶现打不开&#xff0c;报错现象是session not exist&#xff0c;最终定位是redis抖动导致的延迟/断开的现象&#xff0c;最终研发团方案是加入redis…
阅读更多...
排序算法及实现（上）

排序算法及实现（上）

稳定性的判断&#xff1a;如果两个相同大小的元素也进行了交换就是不稳定&#xff0c;否则稳定 1.直接插入排序&#xff1a; 当插入第 i 位置元素时&#xff0c;前面 0 到 i-1 位置的元素已经各自有序。 此时将i 再次从i-1到0位置依次进行比较。找到合适位置将其插入&#x…
阅读更多...
【Python】PYQT5详细介绍

【Python】PYQT5详细介绍

本专栏内容为&#xff1a;Python学习专栏 通过本专栏的深入学习&#xff0c;你可以了解并掌握Python。 &#x1f493;博主csdn个人主页&#xff1a;小小unicorn ⏩专栏分类&#xff1a;Python &#x1f69a;代码仓库&#xff1a;小小unicorn的代码仓库&#x1f69a; &#x1f3…
阅读更多...
GitHub操作

GitHub操作

远程库-GitHub GitHub网址 GitHub是全球最大的远程库 1. 创建远程库 2. 远程仓库操作 2.1 创建远程仓库别名 git remote -v 查看当前所有远程库地址别名 git remote add 别名 远程地址 设置远程库地址别名 案例操作 起一个别名会出现两个别名&#xff0c;是因为既可以拉取…
阅读更多...
macOS DOSBox 汇编环境搭建

macOS DOSBox 汇编环境搭建

正文 一、安装DOSBox 首先前往DOSBox的官网下载并安装最新版本的DOSBox。 二、下载必备的工具包 在用户目录下新建一个文件夹&#xff0c;比如 dosbox: mkdir dosbox然后下载一些常用的工具。下载好了后&#xff0c;将这些工具解压&#xff0c;重新放在 dosbox 这个文件夹…
阅读更多...
数据链路层（详细版）【02】

数据链路层（详细版）【02】

接 数据链路层&#xff08;详细版&#xff09;【01】 文章目录 四、以太网MAC层&#xff08;一&#xff09;MAC地址组成&#xff08;1&#xff09;48位MAC地址格式&#xff08;2&#xff09;单播地址 & 多播地址 & 广播地址&#xff08;3&#xff09;全球管理 & 本…
阅读更多...
笔记2：cifar10数据集获取及pytorch批量处理

笔记2：cifar10数据集获取及pytorch批量处理

&#xff08;1&#xff09;cifar10数据集预处理 CIFAR-10是一个广泛使用的图像数据集&#xff0c;它由10个类别的共60000张32x32彩色图像组成&#xff0c;每个类别有6000张图像。 CIFAR-10官网 以下为CIFAR-10数据集data_batch_*表示训练集数据&#xff0c;test_batch表示测试…
阅读更多...
SpringBoot SpringCloud Vue3 后台权限系统 代码生成器增强版 接私活 毕业设计神器

SpringBoot SpringCloud Vue3 后台权限系统 代码生成器增强版 接私活 毕业设计神器

点击获取项目地址 点击获取项目地址
阅读更多...
Colibri for Mac v2.2.0 原生无损音频播放器 激活版

Colibri for Mac v2.2.0 原生无损音频播放器 激活版

Colibri支持所有流行的无损和有损音频格式的完美清晰的比特完美播放&#xff0c;仅使用微小的计算能力&#xff0c;并提供干净和直观的用户体验。 Colibri在播放音乐时使用极少的计算能力。该应用程序使用最先进的Swift 3编程语言构建&#xff0c;BASS音频引擎作为机器代码捆绑…
阅读更多...
46 udp网络程序

46 udp网络程序

查询网络服务的命令 netstat -nlup n: 显示数字 a&#xff1a;显示所有 u&#xff1a;udp服务 p&#xff1a;显示pid Recv-Q收到的数量&#xff0c;本地ip和远端ip&#xff0c;00表示可以收到任何地址 网络聊天 服务端 定义一个server类&#xff0c;成员保存ip地址&#xff…
阅读更多...
JVM 类加载机制

JVM 类加载机制

JVM 类加载机制分为五个部分&#xff1a;加载&#xff0c;验证&#xff0c;准备&#xff0c;解析&#xff0c;初始化&#xff0c;下面我们就分别来看一下这五个过程。 加载 加载是类加载过程中的一个阶段&#xff0c;这个阶段会在内存中生成一个代表这个类的 java.lang.class 对…
阅读更多...
基于yolov5+streamlit目标检测演示系统设计

基于yolov5+streamlit目标检测演示系统设计

YOLOv5与Streamlit&#xff1a;智能目标检测可视化展示介绍 随着人工智能技术的飞速发展&#xff0c;目标检测技术已成为推动智能化社会进步的关键技术之一。在众多目标检测算法中&#xff0c;YOLOv5以其卓越的性能和实时性&#xff0c;成为了业界的佼佼者。与此同时&#xff…
阅读更多...
UDP多播

UDP多播

1 、多播的概念 多播&#xff0c;也被称为组播&#xff0c;是一种网络通信模式&#xff0c;其中数据的传输和接收仅在同一组内进行。多播具有以下特点&#xff1a; 多播地址标识一组接口&#xff1a;多播使用特定的多播地址&#xff0c;该地址标识一组接收数据的接口。发送到多…
阅读更多...
实现红黑树

实现红黑树

目录 红黑树的概念 红黑树的节点结构定义 红黑树的插入 红黑树的验证 实现红黑树完整代码 红黑树的概念 红黑树 &#xff0c;是一种 二叉搜索树 &#xff0c;但 在每个结点上增加一个存储位表示结点的颜色&#xff0c;可以是 Red 或 Black 。 通过对 任何一条从根到叶子的…
阅读更多...
Leetcode39.组合总和

Leetcode39.组合总和

文章目录 题目描述解题思路重复子集剪枝 代码 题目 参考题解 题目描述 给你一个 无重复元素 的整数数组 candidates 和一个目标整数 target &#xff0c;找出 candidates 中可以使数字和为目标数 target 的 所有 不同组合 &#xff0c;并以列表形式返回。你可以按 任意顺序 返…
阅读更多...
【JavaSE】/*初识Java*/

【JavaSE】/*初识Java*/

目录 一、了解 Java 语言 二、Java 语言的重要性 2.1 使用程度 2.2 工作领域 三、Java 语言的特性 四、Java 的基础语法 五、可能遇到的错误 六、第一个 java 程序代码解析 七、Java 注释 八、Java 标识符 九、Java 关键字 一、了解 Java 语言 Java 是由 Sun Micr…
阅读更多...
最新文章
推荐文章

Copyright @ 2022~2023