[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