llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang @llvm/pr-subscribers-backend-amdgpu Author: Joseph Huber (jhuber6) <details> <summary>Changes</summary> Summary: The backend supports the wavefrontsize intrinsic, and suggests that it is tied to a corresponding clang builtin, but it is not actually present. This simply adds it in so it can be used from clang. This attribute likely isn't the best to rely on, but for the `libc` use-case we will need to detect a struct's differing size in a way that will depend on the wavefront size. --- Full diff: https://github.com/llvm/llvm-project/pull/80741.diff 2 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+1) - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn.cl (+7) ``````````diff diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 5f8001e61a028..213311b96df74 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -69,6 +69,7 @@ BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n") BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n") BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n") BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n") +BUILTIN(__builtin_amdgcn_wavefrontsize, "Ui", "nc") BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZiD*UZiUicC*", "n") BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWiD*UWiUicC*", "n") diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 8d9e4e018b12e..7d9010ee9067d 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -832,6 +832,13 @@ void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) { res = __builtin_amdgcn_atomic_dec32((volatile global uint*)gptr, val, __ATOMIC_SEQ_CST, ""); } +// CHECK-LABEL test_wavefrontsize( +unsigned test_wavefrontsize() { + + // CHECK: call i32 @llvm.amdgcn.wavefrontsize() + return __builtin_amdgcn_wavefrontsize(); +} + // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024} // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025} // CHECK-DAG: attributes #[[$NOUNWIND_READONLY]] = { convergent mustprogress nocallback nofree nounwind willreturn memory(none) } `````````` </details> https://github.com/llvm/llvm-project/pull/80741 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
