add intrinsic clz()

... and fix wrong datatype for CUDA/HIP `ffs()`

Author: psychocoderHPC

include/alpaka/api/syclGeneric/intrinsic.hpp

@@ -57,6 +57,28 @@ namespace alpaka::internal::intrinsic
             ALPAKA_UNREACHABLE(int{});
         }
     };
+
+    template<typename T_Arg>
+    struct Clz::Op<alpaka::internal::SyclIntrinsic, T_Arg>
+    {
+        constexpr auto operator()(alpaka::internal::SyclIntrinsic const, T_Arg const& val) const
+        {
+            if constexpr(sizeof(T_Arg) == 4u)
+            {
+                auto value = std::bit_cast<unsigned int>(val);
+                return sycl::clz(value);
+            }
+            else if constexpr(sizeof(T_Arg) == 8u)
+            {
+                auto value = std::bit_cast<unsigned long long>(val);
+                return sycl::clz(value);
+            }
+            else
+                static_assert(!sizeof(T_Arg), "Unsupported data type, sizeof() must be 4 or 8");
+
+            ALPAKA_UNREACHABLE(int{});
+        }
+    };
 } // namespace alpaka::internal::intrinsic
 
 #endif

include/alpaka/api/unifiedCudaHip/intrinsic.hpp

@@ -43,11 +43,31 @@ namespace alpaka::internal::intrinsic
         {
             if constexpr(sizeof(T_Arg) == 4u)
             {
-                return __ffs(std::bit_cast<unsigned int>(val));
+                return __ffs(std::bit_cast<int>(val));
             }
             else if constexpr(sizeof(T_Arg) == 8u)
             {
-                return __ffsll(std::bit_cast<unsigned long long>(val));
+                return __ffsll(std::bit_cast<long long int>(val));
+            }
+            else
+                static_assert(!sizeof(T_Arg), "Unsupported data type, sizeof() must be 4 or 8");
+
+            ALPAKA_UNREACHABLE(int{});
+        }
+    };
+
+    template<typename T_Arg>
+    struct Clz::Op<alpaka::internal::CudaHipIntrinsic, T_Arg>
+    {
+        inline __device__ auto operator()(alpaka::internal::CudaHipIntrinsic const, T_Arg const& val) const
+        {
+            if constexpr(sizeof(T_Arg) == 4u)
+            {
+                return __clz(std::bit_cast<int>(val));
+            }
+            else if constexpr(sizeof(T_Arg) == 8u)
+            {
+                return __clzll(std::bit_cast<long long int>(val));
             }
             else
                 static_assert(!sizeof(T_Arg), "Unsupported data type, sizeof() must be 4 or 8");

include/alpaka/internal/intrinsic.hpp

@@ -25,4 +25,13 @@ namespace alpaka::internal::intrinsic
             int32_t operator()(T_IntrinsicImpl const, T_Arg const& val) const;
         };
     };
+
+    struct Clz
+    {
+        template<typename T_IntrinsicImpl, typename T_Arg>
+        struct Op
+        {
+            int32_t operator()(T_IntrinsicImpl const, T_Arg const& val) const;
+        };
+    };
 } // namespace alpaka::internal::intrinsic

include/alpaka/internal/stlIntrinsic.hpp

@@ -54,4 +54,24 @@ namespace alpaka::internal::intrinsic
             ALPAKA_UNREACHABLE(int{});
         }
     };
+
+    template<typename T_Arg>
+    struct Clz::Op<alpaka::internal::StlIntrinsic, T_Arg>
+    {
+        constexpr auto operator()(alpaka::internal::StlIntrinsic const, T_Arg const& val) const
+        {
+            if constexpr(sizeof(T_Arg) == 4u)
+            {
+                return std::countl_zero(std::bit_cast<unsigned int>(val));
+            }
+            else if constexpr(sizeof(T_Arg) == 8u)
+            {
+                return std::countl_zero(std::bit_cast<unsigned long long>(val));
+            }
+            else
+                static_assert(!sizeof(T_Arg), "Unsupported data type, sizeof() must be 4 or 8");
+
+            ALPAKA_UNREACHABLE(int{});
+        }
+    };
 } // namespace alpaka::internal::intrinsic

include/alpaka/intrinsic.hpp

@@ -33,4 +33,15 @@ namespace alpaka
         constexpr auto intrinsicImpl = trait::getIntrinsicImpl(thisApi());
         return internal::intrinsic::Ffs::Op<ALPAKA_TYPEOF(intrinsicImpl), ALPAKA_TYPEOF(arg)>{}(intrinsicImpl, arg);
     }
+
+    /* Return the number of most significant zero bits
+     *
+     * @return number consecutive most significant zero bits, zero for input value 0.
+     */
+    constexpr int32_t clz(auto const& arg)
+        requires(sizeof(ALPAKA_TYPEOF(arg)) == 4u || sizeof(ALPAKA_TYPEOF(arg)) == 8u)
+    {
+        constexpr auto intrinsicImpl = alpaka::trait::getIntrinsicImpl(thisApi());
+        return internal::intrinsic::Clz::Op<ALPAKA_TYPEOF(intrinsicImpl), ALPAKA_TYPEOF(arg)>{}(intrinsicImpl, arg);
+    }
 } // namespace alpaka

tests/unit/intrinsic/clz.cpp

@@ -0,0 +1,96 @@
+/*
+ * Copyright 2025 The alpaka team
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <alpaka/alpaka.hpp>
+#include <alpaka/onHost/example/executors.hpp>
+#include <alpaka/onHost/executeForEach.hpp>
+
+#include <catch2/catch_template_test_macros.hpp>
+
+#include <vector>
+
+using namespace alpaka;
+
+using TestBackends
+    = std::decay_t<decltype(onHost::allBackends(onHost::enabledApis, onHost::example::enabledExecutors))>;
+
+struct PopcountKernel
+{
+    template<typename TAcc>
+    ALPAKA_FN_ACC void operator()(TAcc const& acc, concepts::IMdSpan auto output, concepts::IMdSpan auto const input)
+        const
+    {
+        for(auto [index] : onAcc::makeIdxMap(acc, alpaka::onAcc::worker::threadsInGrid, IdxRange{input.getExtents()}))
+        {
+            output[index] = clz(input[index]);
+        }
+    }
+};
+
+TEMPLATE_LIST_TEST_CASE("clz", "[intrinsic][clz]", TestBackends)
+{
+    using Backend = TestType;
+    auto cfg = Backend::makeDict();
+    auto deviceSpec = cfg[object::deviceSpec];
+    auto computeExec = cfg[object::exec];
+
+    // Select a device
+    auto devSelector = alpaka::onHost::makeDeviceSelector(deviceSpec);
+    if(!devSelector.isAvailable())
+    {
+        return;
+    }
+    alpaka::onHost::Device devAcc = devSelector.makeDevice(0);
+
+    // Create a queue on the device
+    alpaka::onHost::Queue queue = devAcc.makeQueue();
+
+    // Input data
+    std::vector<uint64_t> hostInput = {0, 1, 2, 1llu << 32, 1llu << 63};
+    size_t const size = hostInput.size();
+
+    // Allocate device memory
+    auto devInput = alpaka::onHost::alloc<uint64_t>(devAcc, hostInput.size());
+    auto devOutput = alpaka::onHost::alloc<int>(devAcc, hostInput.size());
+
+    // Copy data from host to device
+    alpaka::onHost::memcpy(queue, devInput, hostInput);
+
+    // Define execution parameters
+    auto const frameSpec = alpaka::onHost::getFrameSpec<uint64_t>(devAcc, devInput.getExtents());
+
+    // Create kernel
+    PopcountKernel kernel;
+    auto const taskKernel = alpaka::KernelBundle{kernel, devOutput, devInput};
+
+    // Execute the kernel
+    queue.enqueue(computeExec, frameSpec, taskKernel);
+
+    // Copy data from device to host
+    std::vector<int> hostOutput(size);
+    alpaka::onHost::memcpy(queue, hostOutput, devOutput);
+
+    // Wait for the queue to finish
+    alpaka::onHost::wait(queue);
+
+    // Verification
+    for(size_t i = 0; i < size; ++i)
+    {
+        auto val = hostInput[i];
+        int expected = std::countl_zero(val);
+        CHECK(hostOutput[i] == expected);
+    }
+}