对应第2章 GPU 体系结构上编程模型与 wavefront 执行Task1:对应内容为第 2 章・执行模型核心内容学习时间安排在第 4–6 天建议用时 3 天截止时间为第 6 天结束前当前状态为待安排。学习内容包含三项一是简单了解 GPU 为什么适合并行计算二是认识 thread、block、wavefront不要求掌握底层细节三是结合向量加法示意图理解数组元素如何分给不同线程处理。打卡要求为用几句话或一张草图说明两个数组相加时线程如何分工处理数据。线程如何分工处理数据grid ceil(n/block)个 block每个 block 里block个线程线程用i 块号 × block 块内线程号算出自己负责的元素i ≥ n的线程照样被启动但必须跳过数组读写。此处同样一次运行单元格观察输出即可但是当运行到第五个单元格时会出现以下报错--------------------------------------------------------------------------- CalledProcessError Traceback (most recent call last) File /opt/venv/lib/python3.12/site-packages/torch/utils/cpp_extension.py:2597, in _run_ninja_build(build_directory, verbose, error_prefix) 2596 stdout_fileno 1 - 2597 subprocess.run( 2598 command, 2599 shellIS_WINDOWS and IS_HIP_EXTENSION, 2600 stdoutstdout_fileno if verbose else subprocess.PIPE, 2601 stderrsubprocess.STDOUT, 2602 cwdbuild_directory, 2603 checkTrue, 2604 envenv) 2605 except subprocess.CalledProcessError as e: 2606 # Python 2 and 3 compatible way of getting the error object. File /usr/lib/python3.12/subprocess.py:571, in run(input, capture_output, timeout, check, *popenargs, **kwargs) 570 if check and retcode: -- 571 raise CalledProcessError(retcode, process.args, 572 outputstdout, stderrstderr) 573 return CompletedProcess(process.args, retcode, stdout, stderr) CalledProcessError: Command [ninja, -v] returned non-zero exit status 1. The above exception was the direct cause of the following exception: RuntimeError Traceback (most recent call last) File /workspace/notebooks/common/src/hello_gpu/hip.py:239, in _compile(name, source, preset, entry, grid_limit) 238 with contextlib.redirect_stdout(captured), contextlib.redirect_stderr(captured): -- 239 module load( 240 namemodule_name, 241 sources[str(build_directory / binding.cpp), str(build_directory / kernel.hip)], 242 extra_cflags_FLAGS, 243 extra_cuda_cflags_FLAGS, 244 build_directorystr(build_directory), 245 with_cudaTrue, 246 verboseFalse, 247 ) 248 except Exception as error: File /opt/venv/lib/python3.12/site-packages/torch/utils/cpp_extension.py:1692, in load(name, sources, extra_cflags, extra_cuda_cflags, extra_sycl_cflags, extra_ldflags, extra_include_paths, build_directory, verbose, with_cuda, with_sycl, is_python_module, is_standalone, keep_intermediates) 1597 1598 Load a PyTorch C extension just-in-time (JIT). 1599 (...) 1690 ... verboseTrue) 1691 - 1692 return _jit_compile( 1693 name, 1694 [sources] if isinstance(sources, str) else sources, 1695 extra_cflags, 1696 extra_cuda_cflags, 1697 extra_sycl_cflags, 1698 extra_ldflags, 1699 extra_include_paths, 1700 build_directory or _get_build_directory(name, verbose), 1701 verbose, 1702 with_cuda, 1703 with_sycl, 1704 is_python_module, 1705 is_standalone, 1706 keep_intermediateskeep_intermediates) File /opt/venv/lib/python3.12/site-packages/torch/utils/cpp_extension.py:2134, in _jit_compile(name, sources, extra_cflags, extra_cuda_cflags, extra_sycl_cflags, extra_ldflags, extra_include_paths, build_directory, verbose, with_cuda, with_sycl, is_python_module, is_standalone, keep_intermediates) 2132 sources list(hipified_sources) - 2134 _write_ninja_file_and_build_library( 2135 namename, 2136 sourcessources, 2137 extra_cflagsextra_cflags or [], 2138 extra_cuda_cflagsextra_cuda_cflags or [], 2139 extra_sycl_cflagsextra_sycl_cflags or [], 2140 extra_ldflagsextra_ldflags or [], 2141 extra_include_pathsextra_include_paths or [], 2142 build_directorybuild_directory, 2143 verboseverbose, 2144 with_cudawith_cuda, 2145 with_syclwith_sycl, 2146 is_standaloneis_standalone) 2147 elif verbose: File /opt/venv/lib/python3.12/site-packages/torch/utils/cpp_extension.py:2286, in _write_ninja_file_and_build_library(name, sources, extra_cflags, extra_cuda_cflags, extra_sycl_cflags, extra_ldflags, extra_include_paths, build_directory, verbose, with_cuda, with_sycl, is_standalone) 2285 logger.info(Building extension module %s..., name) - 2286 _run_ninja_build( 2287 build_directory, 2288 verbose, 2289 error_prefixfError building extension {name}) File /opt/venv/lib/python3.12/site-packages/torch/utils/cpp_extension.py:2614, in _run_ninja_build(build_directory, verbose, error_prefix) 2613 message f: {error.output.decode(*SUBPROCESS_DECODE_ARGS)} # type: ignore[union-attr] - 2614 raise RuntimeError(message) from e RuntimeError: Error building extension hello_gpu_65c8a2fdbd3d89cb80e6: [1/3] /opt/rocm/bin/hipcc -DWITH_HIP -DTORCH_EXTENSION_NAMEhello_gpu_65c8a2fdbd3d89cb80e6 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /opt/venv/lib/python3.12/site-packages/torch/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/THH -isystem /opt/rocm/include -isystem /usr/include/python3.12 -fPIC -stdc17 -O3 -stdc17 -D__HIP_PLATFORM_AMD__1 -DUSE_ROCM1 -DHIPBLAS_V2 -fPIC -DCUDA_HAS_FP161 -D__HIP_NO_HALF_OPERATORS__1 -D__HIP_NO_HALF_CONVERSIONS__1 -DHIP_ENABLE_WARP_SYNC_BUILTINS1 --offload-archgfx90a --offload-archgfx942 --offload-archgfx950 --offload-archgfx1100 --offload-archgfx1101 --offload-archgfx1200 --offload-archgfx1201 --offload-archgfx1150 --offload-archgfx1151 -fno-gpu-rdc -O3 -stdc17 -c /workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6/kernel_hip.hip -o kernel_hip.cuda.o FAILED: [code1] kernel_hip.cuda.o /opt/rocm/bin/hipcc -DWITH_HIP -DTORCH_EXTENSION_NAMEhello_gpu_65c8a2fdbd3d89cb80e6 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /opt/venv/lib/python3.12/site-packages/torch/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/THH -isystem /opt/rocm/include -isystem /usr/include/python3.12 -fPIC -stdc17 -O3 -stdc17 -D__HIP_PLATFORM_AMD__1 -DUSE_ROCM1 -DHIPBLAS_V2 -fPIC -DCUDA_HAS_FP161 -D__HIP_NO_HALF_OPERATORS__1 -D__HIP_NO_HALF_CONVERSIONS__1 -DHIP_ENABLE_WARP_SYNC_BUILTINS1 --offload-archgfx90a --offload-archgfx942 --offload-archgfx950 --offload-archgfx1100 --offload-archgfx1101 --offload-archgfx1200 --offload-archgfx1201 --offload-archgfx1150 --offload-archgfx1151 -fno-gpu-rdc -O3 -stdc17 -c /workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6/kernel_hip.hip -o kernel_hip.cuda.o hello_gpu_generated_wrapper:34:21: error: no type named HIPGuardMasqueradingAsCUDA in namespace c10::hip 34 | const c10::hip::HIPGuardMasqueradingAsCUDA guard(a.device()); | ~~~~~~~~~~^ hello_gpu_generated_wrapper:46:21: error: no type named HIPGuardMasqueradingAsCUDA in namespace c10::hip 46 | const c10::hip::HIPGuardMasqueradingAsCUDA guard(a.device()); | ~~~~~~~~~~^ 2 errors generated when compiling for gfx1100. failed to execut[e:/opt/rocm/lib/llvm/bin/clang](file:///E:/opt/rocm/lib/llvm/bin/clang) --offload-archgfx90a --offload-archgfx942 --offload-archgfx950 --offload-archgfx1100 --offload-archgfx1101 --offload-archgfx1200 --offload-archgfx1201 --offload-archgfx1150 --offload-archgfx1151 -DWITH_HIP -DTORCH_EXTENSION_NAMEhello_gpu_65c8a2fdbd3d89cb80e6 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /opt/venv/lib/python3.12/site-packages/torch/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/THH -isystem /opt/rocm/include -isystem /usr/include/python3.12 -fPIC -stdc17 -O3 -stdc17 -D__HIP_PLATFORM_AMD__1 -DUSE_ROCM1 -DHIPBLAS_V2 -fPIC -DCUDA_HAS_FP161 -D__HIP_NO_HALF_OPERATORS__1 -D__HIP_NO_HALF_CONVERSIONS__1 -DHIP_ENABLE_WARP_SYNC_BUILTINS1 -fno-gpu-rdc -O3 -stdc17 -c -x hip /workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6/kernel_hip.hip -o kernel_hip.cuda.o [2/3] c -MMD -MF binding.o.d -DTORCH_EXTENSION_NAMEhello_gpu_65c8a2fdbd3d89cb80e6 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /opt/venv/lib/python3.12/site-packages/torch/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/THH -isystem /opt/rocm/include -isystem /usr/include/python3.12 -fPIC -stdc17 -O3 -stdc17 -c /workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6/binding.cpp -o binding.o -D__HIP_PLATFORM_AMD__1 -DUSE_ROCM1 -DHIPBLAS_V2 -fPIC ninja: build stopped: subcommand failed. The above exception was the direct cause of the following exception: RuntimeError Traceback (most recent call last) Cell In[5], line 1 ---- 1 get_ipython().run_cell_magic(hip, indexed_add --preset binary_f32, \n__global__ void kernel(const float* a, const float* b, float* c, int n) {\n int i blockIdx.x * blockDim.x threadIdx.x;\n if (i n) c[i] a[i] b[i];\n}\n) File /workspace/notebooks/common/src/hello_gpu/hip.py:280, in load_ipython_extension.locals.hip(line, cell) 278 except (SystemExit, ValueError) as error: 279 raise UsageError(Usage: %%hip vector_add [--preset binary_f32] [--entry kernel] [--grid-limit N]) from error -- 280 kernel compile_kernel(arguments.name, cell, presetarguments.preset, 281 entryarguments.entry, grid_limitarguments.grid_limit) 282 ipython.user_ns[arguments.name] kernel 283 action 缓存已复用 if kernel.cache_hit else 编译完成 File /workspace/notebooks/common/src/hello_gpu/hip.py:202, in compile_kernel(name, source, preset, entry, grid_limit) 199 # torchs compiler configuration uses process-wide environment variables. 200 # Serializing compilation also avoids races in its extension-version cache. 201 with _BUILD_LOCK: -- 202 return _compile(name, source, preset, entry, grid_limit) File /workspace/notebooks/common/src/hello_gpu/hip.py:251, in _compile(name, source, preset, entry, grid_limit) 249 log_path build_directory / build.log 250 log_path.write_text(captured.getvalue() \n str(error), encodingutf-8) -- 251 raise RuntimeError( 252 fHIP compilation failed. Kernel lines are labelled notebook_kernel.hip. 253 fFull output: {log_path}\n{error} 254 ) from error 255 finally: 256 if previous_arch is None: RuntimeError: HIP compilation failed. Kernel lines are labelled notebook_kernel.hip. Full output: /workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6/build.log Error building extension hello_gpu_65c8a2fdbd3d89cb80e6: [1/3] /opt/rocm/bin/hipcc -DWITH_HIP -DTORCH_EXTENSION_NAMEhello_gpu_65c8a2fdbd3d89cb80e6 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /opt/venv/lib/python3.12/site-packages/torch/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/THH -isystem /opt/rocm/include -isystem /usr/include/python3.12 -fPIC -stdc17 -O3 -stdc17 -D__HIP_PLATFORM_AMD__1 -DUSE_ROCM1 -DHIPBLAS_V2 -fPIC -DCUDA_HAS_FP161 -D__HIP_NO_HALF_OPERATORS__1 -D__HIP_NO_HALF_CONVERSIONS__1 -DHIP_ENABLE_WARP_SYNC_BUILTINS1 --offload-archgfx90a --offload-archgfx942 --offload-archgfx950 --offload-archgfx1100 --offload-archgfx1101 --offload-archgfx1200 --offload-archgfx1201 --offload-archgfx1150 --offload-archgfx1151 -fno-gpu-rdc -O3 -stdc17 -c /workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6/kernel_hip.hip -o kernel_hip.cuda.o FAILED: [code1] kernel_hip.cuda.o /opt/rocm/bin/hipcc -DWITH_HIP -DTORCH_EXTENSION_NAMEhello_gpu_65c8a2fdbd3d89cb80e6 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /opt/venv/lib/python3.12/site-packages/torch/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/THH -isystem /opt/rocm/include -isystem /usr/include/python3.12 -fPIC -stdc17 -O3 -stdc17 -D__HIP_PLATFORM_AMD__1 -DUSE_ROCM1 -DHIPBLAS_V2 -fPIC -DCUDA_HAS_FP161 -D__HIP_NO_HALF_OPERATORS__1 -D__HIP_NO_HALF_CONVERSIONS__1 -DHIP_ENABLE_WARP_SYNC_BUILTINS1 --offload-archgfx90a --offload-archgfx942 --offload-archgfx950 --offload-archgfx1100 --offload-archgfx1101 --offload-archgfx1200 --offload-archgfx1201 --offload-archgfx1150 --offload-archgfx1151 -fno-gpu-rdc -O3 -stdc17 -c /workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6/kernel_hip.hip -o kernel_hip.cuda.o hello_gpu_generated_wrapper:34:21: error: no type named HIPGuardMasqueradingAsCUDA in namespace c10::hip 34 | const c10::hip::HIPGuardMasqueradingAsCUDA guard(a.device()); | ~~~~~~~~~~^ hello_gpu_generated_wrapper:46:21: error: no type named HIPGuardMasqueradingAsCUDA in namespace c10::hip 46 | const c10::hip::HIPGuardMasqueradingAsCUDA guard(a.device()); | ~~~~~~~~~~^ 2 errors generated when compiling for gfx1100. failed to execut[e:/opt/rocm/lib/llvm/bin/clang](file:///E:/opt/rocm/lib/llvm/bin/clang) --offload-archgfx90a --offload-archgfx942 --offload-archgfx950 --offload-archgfx1100 --offload-archgfx1101 --offload-archgfx1200 --offload-archgfx1201 --offload-archgfx1150 --offload-archgfx1151 -DWITH_HIP -DTORCH_EXTENSION_NAMEhello_gpu_65c8a2fdbd3d89cb80e6 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /opt/venv/lib/python3.12/site-packages/torch/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/THH -isystem /opt/rocm/include -isystem /usr/include/python3.12 -fPIC -stdc17 -O3 -stdc17 -D__HIP_PLATFORM_AMD__1 -DUSE_ROCM1 -DHIPBLAS_V2 -fPIC -DCUDA_HAS_FP161 -D__HIP_NO_HALF_OPERATORS__1 -D__HIP_NO_HALF_CONVERSIONS__1 -DHIP_ENABLE_WARP_SYNC_BUILTINS1 -fno-gpu-rdc -O3 -stdc17 -c -x hip /workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6/kernel_hip.hip -o kernel_hip.cuda.o [2/3] c -MMD -MF binding.o.d -DTORCH_EXTENSION_NAMEhello_gpu_65c8a2fdbd3d89cb80e6 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /opt/venv/lib/python3.12/site-packages/torch/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -isystem /opt/venv/lib/python3.12/site-packages/torch/include/THH -isystem /opt/rocm/include -isystem /usr/include/python3.12 -fPIC -stdc17 -O3 -stdc17 -c /workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6/binding.cpp -o binding.o -D__HIP_PLATFORM_AMD__1 -DUSE_ROCM1 -DHIPBLAS_V2 -fPIC ninja: build stopped: subcommand failed.为了解决这个问题可以新增一个单元格然后粘贴以下内容再运行这个单元格显示运行成功后即可再运行第五个单元格import hello_gpu, pathlib, os, torch, shutil; Gpathlib.Path(hello_gpu.__file__).parent; TG/templates/binary_f32.hip; HG/templates/header.hip; Bpathlib.Path(/workspace/notebooks/.artifacts/build/65c8a2fdbd3d89cb80e6); K(B/kernel.hip).read_text() if (B/kernel.hip).exists() else (目录不在); KH(B/kernel_hip.hip).read_text() if (B/kernel_hip.hip).exists() else (文件不在); sT.read_text(); ns.count(c10::hip::HIPGuardMasqueradingAsCUDA); T.write_text(s.replace(c10::hip::HIPGuardMasqueradingAsCUDA,c10::cuda::CUDAGuard)); LH.read_text().splitlines(); A[] if any(c10/cuda/CUDAGuard.h in x for x in L) else [#include c10/cuda/CUDAGuard.h]; H.write_text(\n.join(LA)\n); print(hello_gpu , hello_gpu.__file__); print(模板换掉旧写法, n, 处 → 现在 旧/新:, T.read_text().count(HIPGuardMasqueradingAsCUDA), /, T.read_text().count(c10::cuda::CUDAGuard), | header 补 include, len(A), 行); print(旧目录 kernel.hip 旧/新:, K.count(HIPGuardMasqueradingAsCUDA), /, K.count(c10::cuda::CUDAGuard), | kernel_hip.hip 旧/新:, KH.count(HIPGuardMasqueradingAsCUDA), /, KH.count(c10::cuda::CUDAGuard)); shutil.rmtree(pathlib.Path(/workspace/notebooks/.artifacts/build), ignore_errorsTrue); os.environ[PYTORCH_ROCM_ARCH]torch.cuda.get_device_properties(0).gcnArchName.split(:)[0]; print(缓存已删 单架构:, os.environ[PYTORCH_ROCM_ARCH])