发表博客之:gemm/threadblock/threadblock_swizzle.h 文件夹讲解,cutlass深入讲解

发表博客之:gemm/threadblock/threadblock_swizzle.h 文件夹讲解,cutlass深入讲解

  • 我们知道,cuda 处理问题都是将一个很大规模的问题分成很多个小问题,每个小问题由一个ThreadBlock来处理,而ThreadblockSwizzle就是负责将逻辑上的小问题映射到cuda上的ThreadBlock上。
  • 或者直接引用这个文件上的注释吧!
  • Implements several possible threadblock-swizzling functions mapping blockIdx to GEMM problems.

先来看一下最简单的struct GemmIdentityThreadblockSwizzle结构体

  • 这个结构体有一个默认参数是1。
template <int N = 1>
struct GemmIdentityThreadblockSwizzle {
   
   

  CUTLASS_HOST_DEVICE
  GemmIdentityThreadblockSwizzle() {
   
    }

  /// Returns the shape of the problem in units of logical tiles
  /// *Gemm* problem size: gemm(M, N, K)
  /// 这个函数的作用是简单的。
  /// 就是以tile_size为逻辑单元,整个问题的逻辑shape!
  CUTLASS_HOST_DEVICE
  static GemmCoord get_tiled_shape(
    GemmCoord problem_size,
    GemmCoord tile_size,
    int split_k_slices) {
   
   

    return GemmCoord(
      (problem_size.m() + tile_size.m() - 1) / tile_size.m(),
      (problem_size.n() + tile_size.n(
class cuda: # CUDA arch to use for CUDA template kernel compilation. # e.g. "70", "75", "80", "90", etc. # When arch is None, Inductor uses torch.cuda.get_device_capability(0). arch: Optional[str] = None # CUDA version to use for CUDA template kernel compilation. # e.g. "11.4", "12.1", etc. # When version is None, Inductor uses torch.version.cuda. version: Optional[str] = None # Optimization level for the host compiler. compile_opt_level: Literal["-O0", "-O1", "-O2", "-O3", "-OS"] = "-O1" # Whether to enable device LTO (link-time-optimization). enable_cuda_lto = False # Whether to keep intermediate files dring compilation. enable_ptxas_info = False # Whether to enable debug info, e.g. line number, cutlass debug info. enable_debug_info = False # Whether to use fast math. use_fast_math = False # Path to the CUTLASS repo root directory. # The default path only works under PyTorch local development environment. cutlass_dir = os.environ.get( "TORCHINDUCTOR_CUTLASS_DIR", os.path.abspath( os.path.join(os.path.dirname(torch.__file__), "../third_party/cutlass/") ), ) # Configures the maximum number of CUTLASS configs to profile in max_autotune. # By default it's None, so that all CUTLASS configs are tuned. # This is mainly used to reduce test time in CI. cutlass_max_profiling_configs: Optional[int] = None # The L2 swizzle values to consider when profiling CUTLASS configs in max_autotune. cutlass_max_profiling_swizzle_options: list[int] = [1, 2, 4] # Path to CUDA NVCC. # NVCC search order: # 1) cuda_cxx set in this config # 2) CUDACXX environment variable # 3) CUDA_HOME environment variable # 4) default system search PATH. cuda_cxx: Optional[str] = None # Minimum value of M*N*K to consider the CUTLASS backend for GEMM ops. cutlass_backend_min_gemm_size: int = 1 # enable generation of inline standalone runner in CUDA CPP generated code # which allows to compile the generated code into a standalone executable. generate_test_runner: bool = ( os.environ.get("INDUCTOR_CUDA_BACKEND_GENERATE_TEST_RUNNER_CODE", "0") == "1" ) # Keep only Cutlass op configs which contain this regular expression pattern # Set this to "warpspecialized_cooperative_epi_tma" to enable only SM90 TMA Cutlass Kernels for large GEMMs cutlass_op_allowlist_regex: Optional[str] = os.environ.get( "TORCHINDUCTOR_CUTLASS_ALLOWLIST" ) # Note: Names of Cutlass ops names can be obtained by calling # op.configuration_name() on a Cutlass op instance, for example those # returned from cutlass_utils.gen_ops() or the op argument passed to # CUTLASSGemmTemplate.render(...) # Filter Cutlass configs which contain this regular expression pattern # Set this to "pingpong" to avoid numerical issues # caused by the op ordering of the "pingpong" memory access # pattern used by some Cutlass Kernels. cutlass_op_denylist_regex: Optional[str] = os.environ.get( "TORCHINDUCTOR_CUTLASS_DENYLIST" ) # Non-negative integer which determines how many kernels are instantiated. # 0 = 0000 generates the fewest kernels, 9999 generates all possible combinations. # increasing first digit reduces schedule / mixed type pruning, # increasing second digit generates more cluster sizes, # increasing third digit generates more MMA multipliers, # increasing fourth digit generates more instruction shapes. cutlass_instantiation_level: str = os.environ.get( "TORCHINDUCTOR_CUTLASS_INSTANTIATION_LEVEL", "0" )
06-12
<think>我们正在解决用户在使用CUDA 12.1和cuDNN编译Darknet时遇到的弃用警告和未声明错误。问题主要涉及两个点: 1. `cudaThreadSynchronize`被弃用(在CUDA 12.1中) 2. `cudnnGetConvolutionForwardAlgorithm`未声明(可能是cuDNN版本变化导致) 根据引用[3]的提示,我们需要注意CUDA的大版本兼容性(如11.*或12.*),但用户已经使用CUDA 12.1,所以我们需要调整代码以适应新版本。 解决方案: 1. **替换弃用的CUDA函数**: - 在CUDA中,`cudaThreadSynchronize`已被弃用,推荐使用`cudaDeviceSynchronize`。我们需要在代码中找到所有`cudaThreadSynchronize`并替换为`cudaDeviceSynchronize`。 2. **更新cuDNN函数调用**: - `cudnnGetConvolutionForwardAlgorithm`在cuDNN v8之后被弃用,改为使用新的API。根据cuDNN的文档,推荐使用`cudnnGetConvolutionForwardAlgorithm_v7`或者使用`cudnnFindConvolutionForwardAlgorithmEx`来获取算法。 - 但是,Darknet项目可能使用的是旧版cuDNN的API,因此我们需要根据cuDNN版本来调整。 具体步骤: 步骤1:替换弃用的CUDA函数 - 在Darknet源码中搜索`cudaThreadSynchronize`,将其替换为`cudaDeviceSynchronize`。 - 注意:这个替换应该在所有出现的地方进行。 步骤2:更新cuDNN相关代码 - 由于`cudnnGetConvolutionForwardAlgorithm`在cuDNN v8中被移除,我们需要根据Darknet中实际使用的情况来修改。 - 在Darknet中,通常是在`src/convolutional_layer.c`或类似文件中调用cuDNN函数。 - 替代方案:使用`cudnnGetConvolutionForwardAlgorithm_v7`(如果仍然可用)或者使用新的API`cudnnFindConvolutionForwardAlgorithmEx`。 但是,由于Darknet可能没有更新这部分代码,我们可以参考以下两种方法: 方法A(推荐):更新到支持cuDNN v8的Darknet分支 - 查看Darknet的GitHub仓库,是否有支持cuDNN v8的分支或提交。如果有,直接更新代码。 方法B:手动修改代码 - 如果不想切换分支,可以手动修改代码。以下是一个修改示例: 原代码(使用旧API): ```c cudnnGetConvolutionForwardAlgorithm(cudnn_handle(), srcTensorDesc, weightDesc, convDesc, dstTensorDesc, CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo); ``` 修改为(使用v7版本的API): ```c int algo_count; cudnnConvolutionFwdAlgoPerf_t algo_perf; cudnnGetConvolutionForwardAlgorithm_v7(cudnn_handle(), srcTensorDesc, weightDesc, convDesc, dstTensorDesc, 1, // 请求的算法数量 &algo_count, &algo_perf); algo = algo_perf.algo; ``` 或者使用`cudnnFindConvolutionForwardAlgorithmEx`(更推荐,因为它可以返回多个算法并选择最优): ```c int algo_count; cudnnConvolutionFwdAlgoPerf_t algo_perf[3]; // 可以测试多个算法 cudnnFindConvolutionForwardAlgorithmEx( cudnn_handle(), srcTensorDesc, input_data, // 输入张量描述和实际数据(需要传入实际数据指针) weightDesc, weight_data, // 权重描述和实际数据 convDesc, dstTensorDesc, output_data, // 输出张量描述和实际数据 3, // 测试的算法数量 &algo_count, algo_perf, workspace, workspace_size); // 需要提供工作空间 // 选择第一个(最优的)算法 algo = algo_perf[0].algo; ``` 注意:使用`cudnnFindConvolutionForwardAlgorithmEx`需要提供工作空间(workspace),这可能需要额外的内存分配。 - 由于Darknet中通常没有提供工作空间,因此修改可能较为复杂。我们可以选择使用`cudnnGetConvolutionForwardAlgorithm_v7`(如果cuDNN版本>=7)来简化。 步骤3:确保cuDNN头文件和库版本匹配 - 检查安装的cuDNN版本是否与CUDA 12.1兼容(cuDNN for CUDA 12.x)。 - 在编译时,确保链接了正确版本的cuDNN库。 步骤4:重新编译 - 修改代码后,执行`make clean`然后重新编译。 注意:由于Darknet项目本身更新较慢,可能对最新CUDA和cuDNN支持不够好。如果上述修改复杂,也可以考虑使用CUDA 11.x和对应cuDNN(根据引用[3]的提示,保持大版本一致,且之前依赖的可能是11.*),这样可能避免代码修改。 但是用户已经使用CUDA 12.1,所以我们提供修改代码的方案。 总结修改步骤: 1. 全局替换`cudaThreadSynchronize`为`cudaDeviceSynchronize`。 2. 修改cuDNN卷积前向算法选择的代码,使用v7版本或FindEx版本的API。 3. 重新编译。 由于Darknet的代码结构可能不同,下面给出一个具体的文件修改指南: - 对于CUDA同步函数的替换: 在命令行中(在darknet源码目录下)执行: ```bash grep -r "cudaThreadSynchronize" . ``` 然后逐个文件修改。 - 对于cuDNN函数,通常在`src/convolutional_layer.c`中的`forward_convolutional_layer_gpu`函数中,找到`cudnnGetConvolutionForwardAlgorithm`调用,并替换为上述新API。 如果用户遇到其他问题,可能需要更详细的调试。 相关问题: 1. 如何检查CUDA和cuDNN的版本? 2. 如何为Darknet指定特定版本的CUDA和cuDNN? 3. 在Darknet中,如何启用或禁用cuDNN加速? 4. 除了上述修改,还有哪些常见的CUDA 12.1编译问题? 5. 如何为Darknet分配工作空间以使用`cudnnFindConvolutionForwardAlgorithmEx`?</think>### 解决CUDA 12.1和cuDNN编译Darknet的问题 以下是针对两个主要问题的解决方案: #### 1. `cudaThreadSynchronize` 弃用警告的解决 **问题原因**: 在CUDA 12.1中,`cudaThreadSynchronize()` 已被弃用,推荐使用 `cudaDeviceSynchronize()`[^3]。 **修复步骤**: 1. 在Darknet源码中全局搜索 `cudaThreadSynchronize` 2. 在以下文件中进行替换: ```c // 修改文件:src/dark_cuda.c // 找到: cudaThreadSynchronize(); // 替换为: cudaDeviceSynchronize(); ``` 3. 同样检查以下文件并替换: - `src/convolutional_layer.c` - `src/deconvolutional_layer.c` - `src/cuda.c` #### 2. `cudnnGetConvolutionForwardAlgorithm` 未声明的解决 **问题原因**: cuDNN v8+ 已弃用旧API,`cudnnGetConvolutionForwardAlgorithm` 被新函数替代[^3]。 **修复方案**: ```c // 修改文件:src/convolutional_layer.c // 找到报错位置(约在forward_convolutional_layer_gpu函数内) // 原始弃用代码: cudnnGetConvolutionForwardAlgorithm(...); // 替换为以下两种方案之一: // 方案A:使用兼容层(推荐) cudnnGetConvolutionForwardAlgorithm_v7( cudnn_handle(), srcTensorDesc, weightDesc, convDesc, dstTensorDesc, 1, // 请求的算法数量 &algo_count, &algo_perf); algo = algo_perf.algo; // 方案B:使用新API cudnnConvolutionFwdAlgoPerf_t perf_results; int algo_count; cudnnFindConvolutionForwardAlgorithmEx( cudnn_handle(), srcTensorDesc, input_data, weightDesc, weight_data, convDesc, dstTensorDesc, output_data, 1, // 测试的算法数量 &algo_count, &perf_results, workspace, workspace_size); // 需要额外分配workspace ``` #### 完整配置建议 1. **环境验证**: ```bash nvcc --version # 确认CUDA 12.1 cat /usr/local/cuda/include/cudnn_version.h | grep CUDNN_MAJOR -A 2 # 确认cuDNN版本 ``` 2. **编译命令**: ```bash make clean # 在Makefile中设置: # GPU=1, CUDNN=1, CUDNN_HALF=1, ARCH=compute_89,sm_89 make -j$(nproc) ``` #### 关键注意事项 1. **cuDNN兼容性**: CUDA 12.1 需要 cuDNN ≥ v8.9.1,检查 `/usr/local/cuda/include/cudnn_version.h` 确认版本[^3] 2. **工作空间分配**: 使用新API时需要额外分配workspace: ```c size_t workspace_size; cudnnGetConvolutionForwardWorkspaceSize(...); cudaMalloc(&workspace, workspace_size); ``` 3. **API差异处理**: 如果遇到其他未定义函数,参考[NVIDIA cuDNN API迁移指南](https://docs.nvidia.com/deeplearning/cudnn/api/index.html#api-migration)进行适配 > **替代方案**:若修改复杂,可降级到CUDA 11.x(需与现有环境一致)[^3]。使用Docker容器管理版本: > ```bash > docker run --gpus all -it nvcr.io/nvidia/cuda:11.8.0-devel-ubuntu20.04 > ``` --- ### 相关问题 1. 如何检查CUDA和cuDNN版本的兼容性? 2. 使用新cuDNN API时如何正确分配工作空间内存? 3. 如何为RTX 40系列显卡优化Darknet性能? 4. 降级CUDA版本时如何避免破坏现有环境? 5. 如何验证修改后的Darknet是否充分利用了GPU? [^1]: 利用CMake编译需要高版本支持,建议使用原生Makefile[^1] [^2]: Windows平台编译需注意VS项目配置[^2] [^3]: CUDA大版本兼容性至关重要,需保持环境一致[^3]
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

fgh431

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值