下载使用Nsight

https://developer.nvidia.com/nsight-systems/get-started

sudo ln -s /opt/nvidia/nsight-systems/2024.4.1/bin/nsys /bin/nsys

nsys profile --stats=true add

API

__ldg

使用内部函数 __ldg() 代替标准指针解引用。__ldg() 强制通过只读数据缓存加载数据，可以有效提高访问效率。

函数实现

1. Sigmoid

mindspore的实现

template <typename T>
__global__ void SigmoidKernel(size_t size, const T *input, T *output) {
  for (int pos = blockIdx.x * blockDim.x + threadIdx.x; pos < size; pos += blockDim.x * gridDim.x) {
    output[pos] = static_cast<T>(1) / (static_cast<T>(1) + exp(-input[pos]));
  }
}

__ldg实现：

#include <cuda.h>
#include <iostream>

template <typename T>
__global__ void SigmoidForwardKernel(const int N, const T* X, T* Y){
    for(size_t i = blockIdx.x * blockDim.x + threadIdx.x; i< N; i+= blockDim.x * gridDim.x){
        #if __CUDA_ARCH__ >= 350
            Y[i] = T(1) / (T(1) + expf(- __ldg(X + i)));
        #else
            Y[i] = T(1) / (T(1) + expf(- X[i]));
        #endif
    }
}

template <typename T>
__global__ void SigmoidBackwardKernel(const int N, const T* dY, const T* Y, T* dX){
    for(size_t i = blockDim.x * blockIdx.x + threadIdx.x; i < N; i += blockDim.x * gridDim.x){
        #if __CUDA_ARCH__ >= 350
            dX[i] = __ldg(dY + i) * __ldg(Y + i) * (T(1) - __ldg(dY + i));
        #else
            dX[i] = dY[i] * Y[i] * (T(1) - Y[i]);
        #endif
    }
}

2. warpReduceSum

// Sums `val` accross all threads in a warp.
//
// Assumptions:
//   - The size of each block should be a multiple of `warpSize`
template <typename T>
__inline__ __device__ T WarpReduceSum(T val) {
#pragma unroll
  for (int offset = (warpSize >> 1); offset > 0; offset >>= 1) {
    val += __shfl_down_sync(0xffffffff, val, offset, warpSize);
  }
  return val;
}

对__shfl_down_sync api的解释参考Warp-Level Primitives。
在mindspore等框架中也有应用，如layer_norm的实现：

template <typename T>
inline __device__ void WarpReduce(T *mean, T *var, T *num) {
  for (int delta = (WARP_SIZE >> 1); delta > 0; delta >>= 1) {
    T mean_other = __shfl_down_sync(0xffffffff, mean[0], delta);
    T var_other = __shfl_down_sync(0xffffffff, var[0], delta);
    T num_other = __shfl_down_sync(0xffffffff, num[0], delta);
    MeanAndVarMerge(mean, var, num, mean_other, var_other, num_other);
  }
}

参考学习资料

• CUDA编程入门及优化
  
• CUDA-Learn-Notes
• 高性能计算与 AI infra
• 如何学习cuda编程？