Merge branch 'next' into stable

.github/workflows/build-cmake.yml

@@ -42,7 +42,7 @@ jobs:
       fail-fast: false
       matrix:
         os: [ macos ]
-        compiler: [ homebrew-clang, system-clang ]
+        compiler: [ homebrew-clang ] # , system-clang
         config: [ Release, Debug ]
     name: ${{ matrix.os }} / ${{ matrix.config }} / ${{ matrix.compiler }}
     runs-on: [ self-hosted, macos, arm64 ]

.github/workflows/build-xmake.yml

@@ -22,15 +22,15 @@ jobs:
         run: |
           xmake lua setup.lua
           xmake f -p linux -a x86_64 --toolchain=${{ matrix.toolchain.name }}-${{ matrix.toolchain.version }} -m ${{ matrix.config }} --cuda_backend=true --enable_dsl=true --enable_gui=true --enable_unity_build=false --enable_tests=true
-          xmake
+          xmake --jobs=16
 
   build-macos:
     strategy:
       fail-fast: false
       matrix:
         os: [ macos ]
-        compiler: [ homebrew-clang, system-clang ]
-        config: [ Release, Debug ]
+        compiler: [ homebrew-clang ] # , system-clang
+        config: [ release, debug ]
     name: ${{ matrix.os }} / ${{ matrix.config }} / ${{ matrix.compiler }}
     runs-on: [ self-hosted, macos, arm64 ]
     steps:
@@ -44,7 +44,7 @@ jobs:
           fi
           xmake lua setup.lua
           xmake f -p macosx --toolchain=llvm -m ${{ matrix.config }} --mm=clang --mxx=clang++ --metal_backend=true --enable_dsl=true --enable_gui=true --enable_unity_build=false --enable_tests=true
-          xmake --jobs=16
+          xmake
 
   build-windows:
     strategy:

.gitignore

@@ -184,6 +184,9 @@ SDKs/
 # Kernel Module Compile Results
 *.mod*
 *.cmd
+*.bat
+*.sh
+*.ps1
 .tmp_versions/
 modules.order
 Module.symvers

include/luisa/ast/attribute.h

@@ -5,6 +5,9 @@ struct Attribute {
     luisa::string key;
     luisa::string value;
     Attribute() noexcept = default;
+    [[nodiscard]] operator bool() const noexcept {
+        return !key.empty();
+    }
     Attribute(
         luisa::string &&key,
         luisa::string &&value) noexcept

include/luisa/ast/type.h

@@ -343,9 +343,9 @@ class LC_AST_API Type {
     /// Return matrix type of type T
     [[nodiscard]] static const Type *matrix(size_t n) noexcept;
     /// Return buffer type of type T
-    [[nodiscard]] static const Type *buffer(const Type *elem) noexcept;
+    [[nodiscard]] static const Type *buffer(const Type *elem, luisa::span<const Attribute> attributes = {}) noexcept;
     /// Return texture type of type T
-    [[nodiscard]] static const Type *texture(const Type *elem, size_t dimension) noexcept;
+    [[nodiscard]] static const Type *texture(const Type *elem, size_t dimension, luisa::span<const Attribute> attributes = {}) noexcept;
     /// Return struct type of type T
     [[nodiscard]] static const Type *structure(luisa::span<Type const *const> members, luisa::span<const Attribute> attributes = {}) noexcept;
     /// Return struct type of type T

include/luisa/core/fiber.h

@@ -95,14 +95,14 @@ template<class F>
     using RetType = decltype(lambda());
     if constexpr (std::is_same_v<RetType, void>) {
         event evt;
-        marl::schedule([evt, lambda = std::forward<F>(lambda)] {
+        marl::schedule([evt, lambda = std::forward<F>(lambda)]() mutable noexcept {
             lambda();
             evt.signal();
         });
         return evt;
     } else {
         future<RetType> evt;
-        marl::schedule([evt, lambda = std::forward<F>(lambda)] {
+        marl::schedule([evt, lambda = std::forward<F>(lambda)]() mutable noexcept {
             evt.signal(lambda());
         });
         return evt;

include/luisa/core/mathematics.h

@@ -52,6 +52,7 @@ using std::fmod;
 using std::round;
 
 using std::exp;
+using std::exp2;
 using std::log;
 using std::log10;
 using std::log2;
@@ -89,6 +90,7 @@ LUISA_MAKE_VECTOR_UNARY_FUNC(floor)
 LUISA_MAKE_VECTOR_UNARY_FUNC(fract)
 LUISA_MAKE_VECTOR_UNARY_FUNC(round)
 LUISA_MAKE_VECTOR_UNARY_FUNC(exp)
+LUISA_MAKE_VECTOR_UNARY_FUNC(exp2)
 LUISA_MAKE_VECTOR_UNARY_FUNC(log)
 LUISA_MAKE_VECTOR_UNARY_FUNC(log10)
 LUISA_MAKE_VECTOR_UNARY_FUNC(log2)

scripts/xmake_func.lua

@@ -8,6 +8,24 @@ set_showmenu(false)
 set_default(false)
 option_end()
 
+option("_lc_check_env")
+set_showmenu(false)
+set_default(false)
+after_check(function(option)
+    if not is_arch("x64", "x86_64", "arm64") then
+        option:set_value(false)
+        utils.error("Illegal environment. Please check your compiler, architecture or platform.")
+        return
+    end
+    if not (is_mode("debug") or is_mode("release") or is_mode("releasedbg")) then
+        option:set_value(false)
+        utils.error("Illegal mode. set mode to 'release', 'debug' or 'releasedbg'.")
+        return
+    end
+    option:set_value(true)
+end)
+option_end()
+
 option("_lc_bin_dir")
 set_default(false)
 set_showmenu(false)
@@ -168,7 +186,7 @@ on_config(function(target)
 end)
 on_load(function(target)
     local _get_or = function(name, default_value)
-        local v = target:values(name)
+        local v = target:extraconf("rules", "lc_basic_settings", name)
         if v == nil then
             return default_value
         end
@@ -231,7 +249,7 @@ on_load(function(target)
     target:add("cxflags", "/Zc:preprocessor", {
         tools = "cl"
     });
-    if _get_or("use_simd", false) then
+    if _get_or("use_simd", get_config("enable_simd")) then
         if is_arch("arm64") then
             target:add("vectorexts", "neon")
         else
@@ -393,26 +411,8 @@ if _disable_unity_build == nil then
         _disable_unity_build = not unity_build
     end
 end
-
-if _configs == nil then
-    _configs = {}
-end
-_configs["use_simd"] = get_config("enable_simd")
 if not _config_project then
     function _config_project(config)
-        if type(_configs) == "table" then
-            for k, v in pairs(_configs) do
-                set_values(k, v)
-            end
-        end
-        if type(_config_rules) == "table" then
-            add_rules(_config_rules)
-        end
-        if type(config) == "table" then
-            for k, v in pairs(config) do
-                set_values(k, v)
-            end
-        end
         local batch_size = config["batch_size"]
         if type(batch_size) == "number" and batch_size > 1 and (not _disable_unity_build) then
             add_rules("c.unity_build", {
@@ -422,5 +422,8 @@ if not _config_project then
                 batchsize = batch_size
             })
         end
+        if type(_config_rules) == "table" then
+            add_rules(_config_rules, config)
+        end
     end
 end

src/ast/type.cpp

@@ -374,6 +374,16 @@ const TypeImpl *TypeRegistry::_decode(luisa::string_view desc) noexcept {
         info->size = (info->size + info->alignment - 1u) / info->alignment * info->alignment;
     } else if (type_identifier == "buffer"sv) {
         info->tag = Type::Tag::BUFFER;
+        while (try_match('[')) {
+            auto attr_key = read_identifier();
+            luisa::string_view attr_value;
+            if (try_match('(')) {
+                attr_value = read_identifier();
+                match(')');
+            }
+            match(']');
+            info->member_attributes.emplace_back(luisa::string{attr_key}, luisa::string{attr_value});
+        }
         match('<');
         auto m = info->members.emplace_back(_decode(split()));
         match('>');
@@ -391,6 +401,16 @@ const TypeImpl *TypeRegistry::_decode(luisa::string_view desc) noexcept {
         info->size = 8u;
     } else if (type_identifier == "texture"sv) {
         info->tag = Type::Tag::TEXTURE;
+        while (try_match('[')) {
+            auto attr_key = read_identifier();
+            luisa::string_view attr_value;
+            if (try_match('(')) {
+                attr_value = read_identifier();
+                match(')');
+            }
+            match(']');
+            info->member_attributes.emplace_back(luisa::string{attr_key}, luisa::string{attr_value});
+        }
         match('<');
         info->dimension = read_number();
         match(',');
@@ -436,8 +456,8 @@ luisa::span<Type const *const> Type::members() const noexcept {
 }
 
 luisa::span<const Attribute> Type::member_attributes() const noexcept {
-    LUISA_ASSERT(is_structure(),
-                 "Calling members() on a non-structure type {}.",
+    LUISA_ASSERT(is_structure() || is_buffer() || is_texture(),
+                 "Calling member_attributes() on a non-structure, buffer or texture type {}.",
                  description());
     return static_cast<const detail::TypeImpl *>(this)->member_attributes;
 }
@@ -586,19 +606,47 @@ const Type *Type::matrix(size_t n) noexcept {
     return from(luisa::format("matrix<{}>", n));
 }
 
-const Type *Type::buffer(const Type *elem) noexcept {
+const Type *Type::buffer(const Type *elem, luisa::span<const Attribute> attributes) noexcept {
     LUISA_ASSERT(!elem->is_buffer() && !elem->is_texture(), "Buffer cannot hold buffers or images.");
     LUISA_ASSERT(!elem->is_structure() || elem->member_attributes().empty(), "Buffer cannot hold structure with custom attributes.");
-    return from(luisa::format("buffer<{}>", elem->description()));
+    if (!attributes.empty()) [[unlikely]] /*usually would not use attribute*/ {
+        luisa::string r{"buffer"};
+        for (auto &attr : attributes) {
+            if (!attr) continue;
+            if (attr.value.empty()) {
+                r.append(luisa::format("[{}]", attr.key));
+            } else {
+                r.append(luisa::format("[{}({})]", attr.key, attr.value));
+            }
+        }
+        r.append(luisa::format("<{}>", elem->description()));
+        return from(r);
+    } else {
+        return from(luisa::format("buffer<{}>", elem->description()));
+    }
 }
 
-const Type *Type::texture(const Type *elem, size_t dimension) noexcept {
+const Type *Type::texture(const Type *elem, size_t dimension, luisa::span<const Attribute> attributes) noexcept {
     if (elem->is_vector()) { elem = elem->element(); }
     LUISA_ASSERT(elem->is_arithmetic(),
                  "Texture element must be an arithmetic, but got {}.",
                  elem->description());
     LUISA_ASSERT(dimension == 2u || dimension == 3u, "Texture dimension must be 2 or 3");
-    return from(luisa::format("texture<{},{}>", dimension, elem->description()));
+    if (!attributes.empty()) [[unlikely]] /*usually would not use attribute*/ {
+        luisa::string r{"texture"};
+        for (auto &attr : attributes) {
+            if (!attr) continue;
+            if (attr.value.empty()) {
+                r.append(luisa::format("[{}]", attr.key));
+            } else {
+                r.append(luisa::format("[{}({})]", attr.key, attr.value));
+            }
+        }
+        r.append(luisa::format("<{},{}>", dimension, elem->description()));
+        return from(r);
+    } else {
+        return from(luisa::format("texture<{},{}>", dimension, elem->description()));
+    }
 }
 
 const Type *Type::structure(size_t alignment, luisa::span<Type const *const> members, luisa::span<const Attribute> attributes) noexcept {
@@ -608,7 +656,7 @@ const Type *Type::structure(size_t alignment, luisa::span<Type const *const> mem
     LUISA_ASSERT(attributes.empty() || attributes.size() == members.size(),
                  "Invalid attribute size (must be empty or same as members' size");
     auto desc = luisa::format("struct<{}", alignment);
-    if (!attributes.empty()) {
+    if (!attributes.empty()) [[unlikely]] /*usually would not use attribute*/ {
         for (size_t i = 0; i < members.size(); ++i) {
             desc.append(",");
             auto &a = attributes[i];

src/backends/common/hlsl/hlsl_codegen_util.cpp

@@ -1856,7 +1856,22 @@ void CodegenUtility::CodegenProperties(
     auto args = kernel.arguments();
     for (auto &&i : vstd::ptr_range(args.data() + offset, args.size() - offset)) {
         auto print = [&] {
-            GetTypeName(*i.type(), varData, kernel.variable_usage(i.uid()));
+            auto usage = kernel.variable_usage(i.uid());
+            if (i.type()->is_buffer() || i.type()->is_texture()) {
+                auto attris = i.type()->member_attributes();
+                if (!attris.empty()) {
+                    for (auto &a : attris) {
+                        if ((to_underlying(usage) & to_underlying(Usage::WRITE)) != 0) {
+                            if (a.key == "cache"sv) {
+                                if (a.value == "coherent"sv) {
+                                    varData << "globallycoherent "sv;
+                                }
+                            }
+                        }
+                    }
+                }
+            }
+            GetTypeName(*i.type(), varData, usage);
             varData << ' ';
             GetVariableName(i, varData);
         };
@@ -1970,15 +1985,28 @@ vstd::MD5 CodegenUtility::GetTypeMD5(vstd::span<Type const *const> types) {
     vstd::vector<uint64_t> typeDescs;
     typeDescs.reserve(types.size());
     for (auto &&i : types) {
-        typeDescs.emplace_back(i->hash());
+        if ((i->is_buffer() || i->is_texture()) && !i->member_attributes().empty())
+            if (i->is_buffer())
+                typeDescs.emplace_back(Type::buffer(i->element())->hash());
+            else
+                typeDescs.emplace_back(Type::texture(i->element(), i->dimension())->hash());
+        else
+            typeDescs.emplace_back(i->hash());
     }
     return {vstd::span<uint8_t const>(reinterpret_cast<uint8_t const *>(typeDescs.data()), typeDescs.size_bytes())};
 }
 vstd::MD5 CodegenUtility::GetTypeMD5(std::initializer_list<vstd::IRange<Variable> *> f) {
     vstd::vector<uint64_t> typeDescs;
     for (auto &&rg : f) {
         for (auto &&i : *rg) {
-            typeDescs.emplace_back(i.type()->hash());
+            auto type = i.type();
+            if ((type->is_buffer() || type->is_texture()) && !type->member_attributes().empty())
+                if (type->is_buffer())
+                    typeDescs.emplace_back(Type::buffer(type->element())->hash());
+                else
+                    typeDescs.emplace_back(Type::texture(type->element(), type->dimension())->hash());
+            else
+                typeDescs.emplace_back(type->hash());
         }
     }
     return {vstd::span<uint8_t const>(reinterpret_cast<uint8_t const *>(typeDescs.data()), typeDescs.size_bytes())};
@@ -1988,7 +2016,14 @@ vstd::MD5 CodegenUtility::GetTypeMD5(Function func) {
     auto args = func.arguments();
     typeDescs.reserve(args.size());
     for (auto &&i : args) {
-        typeDescs.emplace_back(i.type()->hash());
+        auto type = i.type();
+        if ((type->is_buffer() || type->is_texture()) && !type->member_attributes().empty())
+            if (type->is_buffer())
+                typeDescs.emplace_back(Type::buffer(type->element())->hash());
+            else
+                typeDescs.emplace_back(Type::texture(type->element(), type->dimension())->hash());
+        else
+            typeDescs.emplace_back(type->hash());
     }
     return {vstd::span<uint8_t const>(reinterpret_cast<uint8_t const *>(typeDescs.data()), typeDescs.size_bytes())};
 }

src/backends/cuda/cuda_builtin/cuda_device_math.h

@@ -2835,6 +2835,16 @@ template<typename T>
 [[nodiscard]] __device__ inline lc_float3 lc_abs(lc_float3 x) noexcept { return lc_make_float3(fabsf(x.x), fabsf(x.y), fabsf(x.z)); }
 [[nodiscard]] __device__ inline lc_float4 lc_abs(lc_float4 x) noexcept { return lc_make_float4(fabsf(x.x), fabsf(x.y), fabsf(x.z), fabsf(x.w)); }
 
+[[nodiscard]] __device__ inline lc_int lc_abs(lc_int x) noexcept { return abs(x); }
+[[nodiscard]] __device__ inline lc_int2 lc_abs(lc_int2 x) noexcept { return lc_make_int2(abs(x.x), abs(x.y)); }
+[[nodiscard]] __device__ inline lc_int3 lc_abs(lc_int3 x) noexcept { return lc_make_int3(abs(x.x), abs(x.y), abs(x.z)); }
+[[nodiscard]] __device__ inline lc_int4 lc_abs(lc_int4 x) noexcept { return lc_make_int4(abs(x.x), abs(x.y), abs(x.z), abs(x.w)); }
+
+[[nodiscard]] __device__ inline lc_long lc_abs(lc_long x) noexcept { return llabs(x); }
+[[nodiscard]] __device__ inline lc_long2 lc_abs(lc_long2 x) noexcept { return lc_make_long2(llabs(x.x), llabs(x.y)); }
+[[nodiscard]] __device__ inline lc_long3 lc_abs(lc_long3 x) noexcept { return lc_make_long3(llabs(x.x), llabs(x.y), llabs(x.z)); }
+[[nodiscard]] __device__ inline lc_long4 lc_abs(lc_long4 x) noexcept { return lc_make_long4(llabs(x.x), llabs(x.y), llabs(x.z), llabs(x.w)); }
+
 [[nodiscard]] __device__ inline lc_half lc_abs(lc_half x) noexcept { return __habs(x); }
 [[nodiscard]] __device__ inline lc_half2 lc_abs(lc_half2 x) noexcept { return lc_make_half2(__habs(x.x), __habs(x.y)); }
 [[nodiscard]] __device__ inline lc_half3 lc_abs(lc_half3 x) noexcept { return lc_make_half3(__habs(x.x), __habs(x.y), __habs(x.z)); }