[clang] [llvm] [SPIRV] GPU intrinsics (PR #131190)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Mar 13 11:55:56 PDT 2025
github-actions[bot] wrote:
<!--LLVM CODE FORMAT COMMENT: {clang-format}-->
:warning: C/C++ code formatter, clang-format found issues in your code. :warning:
<details>
<summary>
You can test this locally with the following command:
</summary>
``````````bash
git-clang-format --diff 09a36c82793ba32b87faf11fbfc2e7624f25e92e c6eb8e105a711d8433ade0441d37a10d729c70f9 --extensions cpp,c,h -- clang/lib/Headers/spirvintrin.h clang/test/CodeGen/amdgpu-grid-builtins.c clang/test/CodeGen/gpu_builtins.c llvm/include/llvm/Transforms/Scalar/LowerGPUIntrinsic.h llvm/lib/Transforms/Scalar/LowerGPUIntrinsic.cpp clang/lib/Headers/amdgpuintrin.h clang/lib/Headers/gpuintrin.h clang/test/Headers/gpuintrin.c llvm/include/llvm/InitializePasses.h llvm/include/llvm/Transforms/Scalar.h llvm/lib/Passes/PassBuilder.cpp llvm/lib/Transforms/Scalar/Scalar.cpp
``````````
</details>
<details>
<summary>
View the diff from clang-format here.
</summary>
``````````diff
diff --git a/clang/lib/Headers/spirvintrin.h b/clang/lib/Headers/spirvintrin.h
index a5129d3577..1f163bac27 100644
--- a/clang/lib/Headers/spirvintrin.h
+++ b/clang/lib/Headers/spirvintrin.h
@@ -165,7 +165,7 @@ __gpu_match_all_u64(uint64_t __lane_mask, uint64_t __x) {
}
// Terminates execution of the associated wave.
-_DEFAULT_FN_ATTRS [[noreturn]] static __inline__ void __gpu_exit(void) {
+_DEFAULT_FN_ATTRS [[noreturn]] static __inline__ void __gpu_exit(void) {
return __builtin_gpu_exit();
}
diff --git a/llvm/lib/Transforms/Scalar/LowerGPUIntrinsic.cpp b/llvm/lib/Transforms/Scalar/LowerGPUIntrinsic.cpp
index a65e9fa731..6a0b70e52c 100644
--- a/llvm/lib/Transforms/Scalar/LowerGPUIntrinsic.cpp
+++ b/llvm/lib/Transforms/Scalar/LowerGPUIntrinsic.cpp
@@ -412,12 +412,12 @@ static const IntrinsicMap ls[] = {
{gpu_thread_id_x, S<amdgcn_workitem_id_x>, S<nvvm_read_ptx_sreg_tid_x>},
{gpu_thread_id_y, S<amdgcn_workitem_id_y>, S<nvvm_read_ptx_sreg_tid_y>},
{gpu_thread_id_z, S<amdgcn_workitem_id_z>, S<nvvm_read_ptx_sreg_tid_z>},
-
+
// CGBuiltin maps builtin_amdgcn_workgroup_size onto gpu_num_threads
{gpu_num_threads_x, B<amdgpu::WGSize<0>>, S<nvvm_read_ptx_sreg_ntid_x>},
{gpu_num_threads_y, B<amdgpu::WGSize<1>>, S<nvvm_read_ptx_sreg_ntid_y>},
{gpu_num_threads_z, B<amdgpu::WGSize<2>>, S<nvvm_read_ptx_sreg_ntid_z>},
-
+
// Some of the following intrinsics need minor impedance matching
{gpu_num_lanes, S<amdgcn_wavefrontsize>, S<nvvm_read_ptx_sreg_warpsize>},
{gpu_lane_mask, B<amdgpu::lane_mask>, B<nvptx::lane_mask>},
``````````
</details>
https://github.com/llvm/llvm-project/pull/131190
More information about the llvm-commits
mailing list