support cuda device-link

74 files changed, 995 insertions, 815 deletions

diff --git a/.gitignore b/.gitignore
index 33300c4b9..e85cd38da 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,44 +1,29 @@
-*.a
-*.b
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
+# MacOS Cache
+.DS_Store
+
+# Xmake cache
+.xmake/
+build/
+
+# for VS Code
+.vscode/
+
+# for vim
 *.swp
 *.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-*.ilk
-*.stackdump
-*.manifest
-.ccache
-.demo
-.svn
-.DS_Store
-.xmake
-*.deb
-*.zip
-*.gch
-*.gch.d
-*.sys
-cscope.*
-gmon.out
 tags
-doc
-install.log
-build
-core/bin
-core/pre
-core/tool/msys
-core/xmake.config.h
-core/.config.mak
-winenv
-xmake-ppa
-!xmake/actions/build
-.vscode
\ No newline at end of file
+!tags/
+
+# Ignore packaging files
+winenv/
+*.exe
+*.zip
+
+# Makefile generation
+/core/xmake.config.h
+/core/.config.mak
+/core/**/*.o
+/core/**/*.b
+/core/**/*.a
+
+!/xmake/actions/build/
diff --git a/tests/projects/cuda/console/xmake.lua b/tests/projects/cuda/console/xmake.lua
index 430ad7b6d..d7267fb0d 100644
--- a/tests/projects/cuda/console/xmake.lua
+++ b/tests/projects/cuda/console/xmake.lua
@@ -1,6 +1,4 @@
 
-includes('add_cugencodes.lua')
-
 -- define target
 target("cuda_console")
 
diff --git a/tests/projects/cuda/shared/inc/lib.cuh b/tests/projects/cuda/shared/inc/lib.cuh
new file mode 100644
index 000000000..84d5db3ca
--- /dev/null
+++ b/tests/projects/cuda/shared/inc/lib.cuh
@@ -0,0 +1,23 @@
+#pragma once
+
+#include "cuda_runtime.h"
+#include "device_launch_parameters.h"
+
+#ifdef __cplusplus
+extern "C"
+{
+#endif
+
+#if defined(_WIN32)
+#define __export __declspec(dllexport)
+#elif defined(__GNUC__) && ((__GNUC__ >= 4) || (__GNUC__ == 3 && __GNUC_MINOR__ >= 3))
+#define __export __attribute__((visibility("default")))
+#else
+#define __export
+#endif
+
+    __export cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);
+
+#ifdef __cplusplus
+}
+#endif
diff --git a/tests/projects/cuda/shared/src/lib.cu b/tests/projects/cuda/shared/src/lib.cu
new file mode 100644
index 000000000..cd33b167b
--- /dev/null
+++ b/tests/projects/cuda/shared/src/lib.cu
@@ -0,0 +1,97 @@
+#include <lib.cuh>
+#include <stdio.h>
+
+__global__ void addKernel(int *c, const int *a, const int *b)
+{
+    int i = threadIdx.x;
+    c[i] = a[i] + b[i];
+}
+
+// Helper function for using CUDA to add vectors in parallel.
+cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
+{
+    int *dev_a = 0;
+    int *dev_b = 0;
+    int *dev_c = 0;
+    cudaError_t cudaStatus;
+
+    // Choose which GPU to run on, change this on a multi-GPU system.
+    cudaStatus = cudaSetDevice(0);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
+        goto Error;
+    }
+
+    // Allocate GPU buffers for three vectors (two input, one output)    .
+    cudaStatus = cudaMalloc((void **)&dev_c, size * sizeof(int));
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMalloc failed!");
+        goto Error;
+    }
+
+    cudaStatus = cudaMalloc((void **)&dev_a, size * sizeof(int));
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMalloc failed!");
+        goto Error;
+    }
+
+    cudaStatus = cudaMalloc((void **)&dev_b, size * sizeof(int));
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMalloc failed!");
+        goto Error;
+    }
+
+    // Copy input vectors from host memory to GPU buffers.
+    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMemcpy failed!");
+        goto Error;
+    }
+
+    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMemcpy failed!");
+        goto Error;
+    }
+
+    // Launch a kernel on the GPU with one thread for each element.
+    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
+
+    // Check for any errors launching the kernel
+    cudaStatus = cudaGetLastError();
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
+        goto Error;
+    }
+
+    // cudaDeviceSynchronize waits for the kernel to finish, and returns
+    // any errors encountered during the launch.
+    cudaStatus = cudaDeviceSynchronize();
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
+        goto Error;
+    }
+
+    // Copy output vector from GPU buffer to host memory.
+    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMemcpy failed!");
+        goto Error;
+    }
+
+Error:
+    cudaFree(dev_c);
+    cudaFree(dev_a);
+    cudaFree(dev_b);
+
+    return cudaStatus;
+}
diff --git a/tests/projects/cuda/shared/src/main.cu b/tests/projects/cuda/shared/src/main.cu
new file mode 100644
index 000000000..a33ec5c3e
--- /dev/null
+++ b/tests/projects/cuda/shared/src/main.cu
@@ -0,0 +1,34 @@
+
+#include "cuda_runtime.h"
+#include <stdio.h>
+#include <lib.cuh>
+
+int main()
+{
+    const int arraySize = 5;
+    const int a[arraySize] = {1, 2, 3, 4, 5};
+    const int b[arraySize] = {10, 20, 30, 40, 50};
+    int c[arraySize] = {0};
+
+    // Add vectors in parallel.
+    cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "addWithCuda failed!");
+        return 1;
+    }
+
+    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
+           c[0], c[1], c[2], c[3], c[4]);
+
+    // cudaDeviceReset must be called before exiting in order for profiling and
+    // tracing tools such as Nsight and Visual Profiler to show complete traces.
+    cudaStatus = cudaDeviceReset();
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaDeviceReset failed!");
+        return 1;
+    }
+
+    return 0;
+}
diff --git a/tests/projects/cuda/shared/xmake.lua b/tests/projects/cuda/shared/xmake.lua
new file mode 100644
index 000000000..e96999a13
--- /dev/null
+++ b/tests/projects/cuda/shared/xmake.lua
@@ -0,0 +1,22 @@
+
+-- add modes: debug and release
+add_rules("mode.debug", "mode.release")
+
+-- generate PTX code for the virtual architecture to guarantee compatibility
+add_cugencodes("compute_30")
+
+-- define target
+target("lib")
+
+    -- set kind
+    set_kind("shared")
+
+    -- add files
+    add_files("src/lib.cu")
+
+    add_includedirs("inc", {public = true})
+
+target("bin")
+    add_deps("lib")
+    set_kind("binary")
+    add_files("src/main.cu")
diff --git a/tests/projects/cuda/static/inc/lib.cuh b/tests/projects/cuda/static/inc/lib.cuh
new file mode 100644
index 000000000..35255e31e
--- /dev/null
+++ b/tests/projects/cuda/static/inc/lib.cuh
@@ -0,0 +1,3 @@
+#pragma once
+
+__global__ void addKernel(int *c, const int *a, const int *b);
\ No newline at end of file
diff --git a/tests/projects/cuda/static/src/lib.cu b/tests/projects/cuda/static/src/lib.cu
new file mode 100644
index 000000000..a5f054354
--- /dev/null
+++ b/tests/projects/cuda/static/src/lib.cu
@@ -0,0 +1,7 @@
+#include <lib.cuh>
+
+__global__ void addKernel(int *c, const int *a, const int *b)
+{
+    int i = threadIdx.x;
+    c[i] = a[i] + b[i];
+}
\ No newline at end of file
diff --git a/tests/projects/cuda/static/src/main.cu b/tests/projects/cuda/static/src/main.cu
new file mode 100644
index 000000000..32844cd56
--- /dev/null
+++ b/tests/projects/cuda/static/src/main.cu
@@ -0,0 +1,127 @@
+
+#include "cuda_runtime.h"
+#include "device_launch_parameters.h"
+
+#include <stdio.h>
+#include <lib.cuh>
+
+cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);
+
+int main()
+{
+    const int arraySize = 5;
+    const int a[arraySize] = {1, 2, 3, 4, 5};
+    const int b[arraySize] = {10, 20, 30, 40, 50};
+    int c[arraySize] = {0};
+
+    // Add vectors in parallel.
+    cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "addWithCuda failed!");
+        return 1;
+    }
+
+    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
+           c[0], c[1], c[2], c[3], c[4]);
+
+    // cudaDeviceReset must be called before exiting in order for profiling and
+    // tracing tools such as Nsight and Visual Profiler to show complete traces.
+    cudaStatus = cudaDeviceReset();
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaDeviceReset failed!");
+        return 1;
+    }
+
+    return 0;
+}
+
+// Helper function for using CUDA to add vectors in parallel.
+cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
+{
+    int *dev_a = 0;
+    int *dev_b = 0;
+    int *dev_c = 0;
+    cudaError_t cudaStatus;
+
+    // Choose which GPU to run on, change this on a multi-GPU system.
+    cudaStatus = cudaSetDevice(0);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
+        goto Error;
+    }
+
+    // Allocate GPU buffers for three vectors (two input, one output)    .
+    cudaStatus = cudaMalloc((void **)&dev_c, size * sizeof(int));
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMalloc failed!");
+        goto Error;
+    }
+
+    cudaStatus = cudaMalloc((void **)&dev_a, size * sizeof(int));
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMalloc failed!");
+        goto Error;
+    }
+
+    cudaStatus = cudaMalloc((void **)&dev_b, size * sizeof(int));
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMalloc failed!");
+        goto Error;
+    }
+
+    // Copy input vectors from host memory to GPU buffers.
+    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMemcpy failed!");
+        goto Error;
+    }
+
+    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMemcpy failed!");
+        goto Error;
+    }
+
+    // Launch a kernel on the GPU with one thread for each element.
+    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
+
+    // Check for any errors launching the kernel
+    cudaStatus = cudaGetLastError();
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
+        goto Error;
+    }
+
+    // cudaDeviceSynchronize waits for the kernel to finish, and returns
+    // any errors encountered during the launch.
+    cudaStatus = cudaDeviceSynchronize();
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
+        goto Error;
+    }
+
+    // Copy output vector from GPU buffer to host memory.
+    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMemcpy failed!");
+        goto Error;
+    }
+
+Error:
+    cudaFree(dev_c);
+    cudaFree(dev_a);
+    cudaFree(dev_b);
+
+    return cudaStatus;
+}
diff --git a/tests/projects/cuda/static/xmake.lua b/tests/projects/cuda/static/xmake.lua
new file mode 100644
index 000000000..55e1e5fd9
--- /dev/null
+++ b/tests/projects/cuda/static/xmake.lua
@@ -0,0 +1,24 @@
+
+-- add modes: debug and release
+add_rules("mode.debug", "mode.release")
+
+-- generate PTX code for the virtual architecture to guarantee compatibility
+add_cugencodes("compute_30")
+
+-- define target
+target("lib")
+
+    -- set kind
+    set_kind("static")
+
+    add_cuflags("-rdc=true")
+
+    add_includedirs("inc", {public = true})
+
+    -- add files
+    add_files("src/lib.cu")
+
+target("bin")
+    add_deps("lib")
+    set_kind("binary")
+    add_files("src/main.cu")
diff --git a/xmake/core/language/language.lua b/xmake/core/language/language.lua
index a7fdddcc0..ee57172fe 100644
--- a/xmake/core/language/language.lua
+++ b/xmake/core/language/language.lua
@@ -103,6 +103,11 @@ function _instance:extensions()
     return extensions
 end
 
+-- get the rules
+function _instance:rules()
+    return self._INFO:get("rules")
+end
+
 -- get the source kinds
 function _instance:sourcekinds()
     return self._INFO:get("sourcekinds")
@@ -221,6 +226,8 @@ function language._interpreter()
         {
             -- language.set_xxx
             "language.set_mixingkinds"
+            -- language.add_xxx
+        ,   "language.add_rules"
         }
     ,   script =
         {
diff --git a/xmake/core/project/deprecated/project.lua b/xmake/core/project/deprecated/project.lua
index b55970c86..85d05ed8f 100644
--- a/xmake/core/project/deprecated/project.lua
+++ b/xmake/core/project/deprecated/project.lua
@@ -27,6 +27,7 @@ local path                      = require("base/path")
 local utils                     = require("base/utils")
 local table                     = require("base/table")
 local string                    = require("base/string")
+local rule                      = require("project/rule")
 local config                    = require("project/config")
 local platform                  = require("platform/platform")
 local deprecated                = require("base/deprecated")
diff --git a/xmake/core/project/project.lua b/xmake/core/project/project.lua
index 07fb6792f..6e670ec3b 100644
--- a/xmake/core/project/project.lua
+++ b/xmake/core/project/project.lua
@@ -498,7 +498,7 @@ function project._load_targets()
         t._ORDERDEPS = t._ORDERDEPS or {}
         project._load_deps(t, targets, t._DEPS, t._ORDERDEPS)
 
-        -- load rules
+        -- load rules from target and language
         --
         -- .e.g 
         --
@@ -509,7 +509,15 @@ function project._load_targets()
         --
         t._RULES      = t._RULES or {}
         t._ORDERULES  = t._ORDERULES or {}
-        for _, rulename in ipairs(table.wrap(t:get("rules"))) do
+        local rulenames = {}
+        table.join2(rulenames, t:get("rules"))
+        for _, sourcefile in ipairs(table.wrap(t:get("files"))) do
+            local lang = language.load_ex(path.extension(sourcefile))
+            if lang and lang:rules() then
+                table.join2(rulenames, lang:rules())
+            end
+        end
+        for _, rulename in ipairs(rulenames) do
             local r = project.rule(rulename) or rule.rule(rulename)
             if r then
                 t._RULES[rulename] = r
diff --git a/xmake/core/project/template.lua b/xmake/core/project/template.lua
index 8f6e959f1..27dc34e42 100644
--- a/xmake/core/project/template.lua
+++ b/xmake/core/project/template.lua
@@ -280,7 +280,7 @@ function template.create(language, templateid, targetname)
     -- append FAQ to xmake.lua
     local projectfile = path.join(projectdir, "xmake.lua")
     if os.isfile(projectfile) then
-        local file = io.open("xmake.lua", "a+")
+        local file = io.open(projectfile, "a+")
         if file then
             file:print("")
             file:print(template.faq())
@@ -288,6 +288,9 @@ function template.create(language, templateid, targetname)
         end
     end
 
+    -- generate .gitignore
+    os.cp(path.join(os.programdir(), "scripts", "gitignore"), path.join(projectdir, ".gitignore"))
+
     -- ok
     return true
 end
diff --git a/xmake/core/sandbox/modules/os.lua b/xmake/core/sandbox/modules/os.lua
index c7ae6d8d3..2fc508fff 100644
--- a/xmake/core/sandbox/modules/os.lua
+++ b/xmake/core/sandbox/modules/os.lua
@@ -327,7 +327,7 @@ function sandbox_os.vrunv(program, argv, opt)
 
     -- echo command
     if option.get("verbose") then
-        print(vformat(program), table.concat(argv, " "))
+        print(vformat(program) .. " " .. table.concat(argv, " "))
     end
 
     -- run it
diff --git a/xmake/core/tool/linker.lua b/xmake/core/tool/linker.lua
index 78d6a85f1..a82b12918 100644
--- a/xmake/core/tool/linker.lua
+++ b/xmake/core/tool/linker.lua
@@ -132,7 +132,7 @@ function linker.load(targetkind, sourcekinds, target)
     local linkerinfo = linkerinfo_or_errors
 
     -- init cache key
-    local cachekey = linkerinfo.linkerkind .. (linkerinfo.program or "") .. (config.get("arch") or os.arch())
+    local cachekey = targetkind .. "_" .. linkerinfo.linkerkind .. (linkerinfo.program or "") .. (config.get("arch") or os.arch())
 
     -- get it directly from cache dirst
     builder._INSTANCES = builder._INSTANCES or {}
@@ -242,12 +242,6 @@ function linker:linkflags(opt)
     -- add flags from the platform 
     self:_add_flags_from_platform(flags, targetkind)
 
-    --[[
-    -- add flags from the compiler 
-    if target then
-        self:_add_flags_from_compiler(flags, target, targetkind)
-    end]]
-
     -- add flags from the linker 
     self:_add_flags_from_linker(flags)
 
diff --git a/xmake/languages/cuda/api.lua b/xmake/languages/cuda/api.lua
index 589eb5579..f825c9b72 100644
--- a/xmake/languages/cuda/api.lua
+++ b/xmake/languages/cuda/api.lua
@@ -27,7 +27,9 @@ function apis()
         -- target.add_xxx
         "target.add_links"
     ,   "target.add_syslinks"
+    ,   "target.add_cugencodes"
     ,   "target.add_cuflags"
+    ,   "target.add_culdflags"
     ,   "target.add_ldflags"
     ,   "target.add_arflags"
     ,   "target.add_shflags"
@@ -37,7 +39,9 @@ function apis()
         -- option.add_xxx
     ,   "option.add_links"
     ,   "option.add_syslinks"
+    ,   "option.add_cugencodes"
     ,   "option.add_cuflags"
+    ,   "option.add_culdflags"
     ,   "option.add_ldflags"
     ,   "option.add_arflags"
     ,   "option.add_shflags"
diff --git a/xmake/languages/cuda/xmake.lua b/xmake/languages/cuda/xmake.lua
index 46ba8f902..75b4bf5b8 100644
--- a/xmake/languages/cuda/xmake.lua
+++ b/xmake/languages/cuda/xmake.lua
@@ -28,10 +28,10 @@ language("cuda")
     set_sourceflags {cu = "cuflags"}
 
     -- set target kinds
-    set_targetkinds {binary = "cu-ld", static = "cu-ar", shared = "cu-sh"}
+    set_targetkinds {gpucode = "cu-ld", binary = "ld", static = "ar", shared = "sh"}
 
     -- set target flags
-    set_targetflags {binary = "ldflags", static = "arflags", shared = "shflags"}
+    set_targetflags {gpucode = "culdflags", binary = "ldflags", static = "arflags", shared = "shflags"}
 
     -- set language kinds
     set_langkinds {cu = "cu"}
@@ -39,6 +39,9 @@ language("cuda")
     -- set mixing kinds
     set_mixingkinds("cu", "cc", "cxx", "as")
 
+    -- add rules
+    add_rules("cuda")
+
     -- on load
     on_load("load")
 
@@ -113,6 +116,21 @@ language("cuda")
             "target.strip"
         ,   "target.symbols"
         }
+    ,   gpucode = 
+        {
+            "config.linkdirs"
+        ,   "target.linkdirs"
+        ,   "option.linkdirs"
+        ,   "platform.linkdirs"
+        ,   "config.links"
+        ,   "target.links"
+        ,   "option.links"
+        ,   "platform.links"
+        ,   "config.syslinks"
+        ,   "target.syslinks"
+        ,   "option.syslinks"
+        ,   "platform.syslinks"
+        }
     }
 
     -- set menu
@@ -121,7 +139,7 @@ language("cuda")
                 {   
                     {category = "Cross Complation Configuration/Compiler Configuration"         }
                 ,   {nil, "cu",         "kv", nil,          "The Cuda Compiler"                 }
-                ,   {nil, "cu-cxx",     "kv", nil,          "The Cuda Host C++ Compiler"        }
+                ,   {nil, "cu-ccbin",   "kv", nil,          "The Cuda Host C++ Compiler"        }
 
                 ,   {category = "Cross Complation Configuration/Linker Configuration"           }
                 ,   {nil, "cu-ld",      "kv", nil,          "The Cuda Linker"                   }
diff --git a/xmake/modules/core/tools/cl.lua b/xmake/modules/core/tools/cl.lua
index 3859358d3..4f3545988 100644
--- a/xmake/modules/core/tools/cl.lua
+++ b/xmake/modules/core/tools/cl.lua
@@ -313,10 +313,10 @@ function _include_deps(self, outdata)
     -- translate it
     local results = {}
     local uniques = {}
-    for _, line in ipairs(outdata:split("\r\n", {plain = true})) do
+    for _, line in ipairs(outdata:split("\n", {plain = true})) do
 
         -- get includefile
-        local includefile = _include_note(self, line)
+        local includefile = _include_note(self, line:trim())
         if includefile then
 
             -- get the relative
diff --git a/xmake/modules/core/tools/nvcc.lua b/xmake/modules/core/tools/nvcc.lua
index 568c6aef9..592634ac8 100644
--- a/xmake/modules/core/tools/nvcc.lua
+++ b/xmake/modules/core/tools/nvcc.lua
@@ -29,9 +29,6 @@ import("private.tools.nvcc.parse_deps")
 -- init it
 function init(self)
 
-    -- init shflags
-    self:set("cu-shflags", "-shared")
-
     -- init flags
     if not is_plat("windows") then
         self:set("shared.cuflags", "-Xcompiler -fPIC")
diff --git a/xmake/modules/detect/sdks/find_cuda.lua b/xmake/modules/detect/sdks/find_cuda.lua
index 259f643fe..f9077a485 100644
--- a/xmake/modules/detect/sdks/find_cuda.lua
+++ b/xmake/modules/detect/sdks/find_cuda.lua
@@ -19,10 +19,14 @@
 --
 
 -- imports
+import("lib.detect.cache")
 import("lib.detect.find_file")
+import("core.base.option")
+import("core.base.global")
+import("core.project.config")
 
 -- find cuda sdk directory
-function _find_cudadir()
+function _find_sdkdir()
 
     -- init the search directories
     local pathes = {}
@@ -42,11 +46,50 @@ function _find_cudadir()
 end
 
 -- find cuda sdk toolchains
+function _find_cuda(sdkdir)
+
+    -- find cuda directory
+    if not sdkdir or not os.isdir(sdkdir) then
+        sdkdir = _find_sdkdir()
+    end
+
+    -- not found?
+    if not sdkdir or not os.isdir(sdkdir) then
+        return nil
+    end
+
+    -- get the bin directory 
+    local bindir = path.join(sdkdir, "bin")
+    if not os.isexec(path.join(bindir, "nvcc")) then
+        return nil
+    end
+
+    -- get linkdirs
+    local linkdirs = {}
+    if is_plat("windows") then
+        local subdir = is_arch("x64") and "x64" or "Win32"
+        table.insert(linkdirs, path.join(sdkdir, "lib", subdir))
+    elseif is_plat("linux") and is_arch("x86_64") then
+        table.insert(linkdirs, path.join(sdkdir, "lib64", "stubs"))
+        table.insert(linkdirs, path.join(sdkdir, "lib64"))
+    else
+        table.insert(linkdirs, path.join(sdkdir, "lib", "stubs"))
+        table.insert(linkdirs, path.join(sdkdir, "lib"))
+    end
+
+    -- get includedirs
+    local includedirs = {path.join(sdkdir, "include")}
+
+    -- get toolchains
+    return {sdkdir = sdkdir, bindir = bindir, linkdirs = linkdirs, includedirs = includedirs}
+end
+
+-- find cuda sdk toolchains
 --
--- @param cudadir   the cuda directory
+-- @param sdkdir    the cuda sdk directory
 -- @param opt       the argument options 
 --
--- @return          the cuda sdk toolchains. .e.g {cudadir = ..., bindir = .., linkdirs = ..., includedirs = ..., .. }
+-- @return          the cuda sdk toolchains. .e.g {sdkdir = ..., bindir = .., linkdirs = ..., includedirs = ..., .. }
 --
 -- @code 
 --
@@ -54,33 +97,41 @@ end
 -- 
 -- @endcode
 --
-function main(cudadir, opt)
+function main(sdkdir, opt)
 
     -- init arguments
     opt = opt or {}
 
-    -- find cuda directory
-    if not cudadir or not os.isdir(cudadir) then
-        cudadir = _find_cudadir()
+    -- attempt to load cache first
+    local key = "detect.sdks.find_cuda"
+    local cacheinfo = cache.load(key)
+    if not opt.force and cacheinfo.cuda and cacheinfo.cuda.sdkdir and os.isdir(cacheinfo.cuda.sdkdir) then
+        return cacheinfo.cuda
     end
+       
+    -- find cuda
+    local cuda = _find_cuda(sdkdir or config.get("cuda") or global.get("cuda") or config.get("sdk"))
+    if cuda then
 
-    -- not found?
-    if not cudadir or not os.isdir(cudadir) then
-        return nil
-    end
+        -- save to config
+        config.set("cuda", cuda.sdkdir, {force = true, readonly = true})
 
-    -- get the bin directory 
-    local bindir = path.join(cudadir, "bin")
-    if not os.isexec(path.join(bindir, "nvcc")) then
-        return nil
-    end
+        -- trace
+        if opt.verbose or option.get("verbose") then
+            cprint("checking for the Cuda SDK directory ... ${color.success}%s", cuda.sdkdir)
+        end
+    else
 
-    -- get linkdirs
-    local linkdirs = {path.join(cudadir, "lib")}
+        -- trace
+        if opt.verbose or option.get("verbose") then
+            cprint("checking for the Cuda SDK directory ... ${color.nothing}${text.nothing}")
+        end
+    end
 
-    -- get includedirs
-    local includedirs = {path.join(cudadir, "include")}
+    -- save to cache
+    cacheinfo.cuda = cuda or false
+    cache.save(key, cacheinfo)
 
-    -- get toolchains
-    return {cudadir = cudadir, bindir = bindir, linkdirs = linkdirs, includedirs = includedirs}
+    -- ok?
+    return cuda
 end
diff --git a/xmake/modules/detect/sdks/find_qt.lua b/xmake/modules/detect/sdks/find_qt.lua
index 885198575..6ab39b254 100644
--- a/xmake/modules/detect/sdks/find_qt.lua
+++ b/xmake/modules/detect/sdks/find_qt.lua
@@ -153,9 +153,9 @@ function main(sdkdir, opt)
     opt = opt or {}
 
     -- attempt to load cache first
-    local key = "detect.sdks.find_qt." .. (sdkdir or "")
+    local key = "detect.sdks.find_qt"
     local cacheinfo = cache.load(key)
-    if not opt.force and cacheinfo.qt then
+    if not opt.force and cacheinfo.qt and cacheinfo.qt.sdkdir and os.isdir(cacheinfo.qt.sdkdir) then
         return cacheinfo.qt
     end
        
diff --git a/xmake/modules/detect/sdks/find_wdk.lua b/xmake/modules/detect/sdks/find_wdk.lua
index 513be1f88..e19befd02 100644
--- a/xmake/modules/detect/sdks/find_wdk.lua
+++ b/xmake/modules/detect/sdks/find_wdk.lua
@@ -160,9 +160,9 @@ function main(sdkdir, opt)
     opt = opt or {}
 
     -- attempt to load cache first
-    local key = "detect.sdks.find_wdk." .. (sdkdir or "")
+    local key = "detect.sdks.find_wdk"
     local cacheinfo = cache.load(key)
-    if not opt.force and cacheinfo.wdk then
+    if not opt.force and cacheinfo.wdk and cacheinfo.wdk.sdkdir and os.isdir(cacheinfo.wdk.sdkdir) then
         return cacheinfo.wdk
     end
        
diff --git a/xmake/modules/lib/detect/find_cudadevices.lua b/xmake/modules/lib/detect/find_cudadevices.lua
index 64a69f705..aa44fe0fb 100644
--- a/xmake/modules/lib/detect/find_cudadevices.lua
+++ b/xmake/modules/lib/detect/find_cudadevices.lua
@@ -21,6 +21,7 @@
 -- imports
 import("core.base.option")
 import("core.platform.platform")
+import("core.project.config")
 import("lib.detect.cache")
 
 -- a magic string to filter output
@@ -120,13 +121,15 @@ function _find_devices(verbose)
         cprint("${dim}checking for cuda devices")
     end
 
-    local sourcefile = path.join(os.programdir(), 'scripts', 'find_cudadevices.cu')
+    local sourcefile = path.join(os.programdir(), 'scripts', 'find_cudadevices.cpp')
     local outfile = os.tmpfile()
     local compile_errors = nil
     local results, errors = try 
     { 
-        function () 
-            return os.iorunv(nvcc, { sourcefile, '-run', '-o', outfile , '-DPRINT_SUFFIX="' .. _PRINT_SUFFIX .. '"' }) 
+        function ()
+            local archs = { i386 = "-m32", x86 = "-m32", x86_64 = "-m64", x64 = "-m64" }
+            local arch = archs[config.get("arch")] or ""
+            return os.iorunv(nvcc, { sourcefile, arch, '-run', '-o', outfile , '-DPRINT_SUFFIX="' .. _PRINT_SUFFIX .. '"' }) 
         end, 
         catch 
         {
diff --git a/xmake/modules/private/platform/check_cuda.lua b/xmake/modules/private/platform/check_cuda.lua
index 0ab59a28b..ba86fbb01 100644
--- a/xmake/modules/private/platform/check_cuda.lua
+++ b/xmake/modules/private/platform/check_cuda.lua
@@ -24,21 +24,9 @@ import("detect.sdks.find_cuda")
 
 -- check the cuda sdk toolchains
 function main(config)
-
-    -- get the cuda directory
-    local cuda_dir = config.get("cuda")
-    if not cuda_dir then
-
-        -- check ok? update it
-        local toolchains = find_cuda()
-        if toolchains then
-
-            -- save it
-            config.set("cuda", toolchains.cudadir)
-
-            -- trace
-            cprint("checking for the Cuda SDK directory ... ${green}%s", toolchains.cudadir)
-        end
+    local cuda = find_cuda(config.get("cuda"), {verbose = true})
+    if cuda then
+        config.set("cuda", cuda.sdkdir, {force = true, readonly = true})
     end
 end
 
diff --git a/xmake/platforms/linux/config.lua b/xmake/platforms/linux/config.lua
index 60cc98676..4cbed376d 100644
--- a/xmake/platforms/linux/config.lua
+++ b/xmake/platforms/linux/config.lua
@@ -78,13 +78,13 @@ function _toolchains()
     local cu         = toolchain("the cuda compiler")
     local cu_ld      = toolchain("the cuda linker")
     local cu_sh      = toolchain("the cuda shared library linker")
-    local cu_cxx     = toolchain("the cuda host c++ compiler")
+    local cu_ccbin     = toolchain("the cuda host c++ compiler")
     local toolchains = {cc = cc, cxx = cxx, as = as, ld = ld, sh = sh, ar = ar, ex = ex, 
                         mm = mm, mxx = mxx,
                         gc = gc, ["gc-ld"] = gc_ld, ["gc-ar"] = gc_ar,
                         dc = dc, ["dc-ld"] = dc_ld, ["dc-sh"] = dc_sh, ["dc-ar"] = dc_ar,
                         rc = rc, ["rc-ld"] = rc_ld, ["rc-sh"] = rc_sh, ["rc-ar"] = rc_ar,
-                        cu = cu, ["cu-ld"] = cu_ld, ["cu-sh"] = cu_sh, ["cu-cxx"] = cu_cxx}
+                        cu = cu, ["cu-ld"] = cu_ld, ["cu-sh"] = cu_sh, ["cu-ccbin"] = cu_ccbin}
 
     -- init the c compiler
     cc:add("$(env CC)", {name = "gcc", cross = cross}, {name = "clang", cross = cross})
@@ -151,7 +151,7 @@ function _toolchains()
     cu_ld:add("nvcc")
     cu_sh:add("nvcc")
     if not cross or cross == "" then
-        cu_cxx:add("$(env CXX)", "$(env CC)", "gcc", "clang", "g++", "clang++")
+        cu_ccbin:add("$(env CXX)", "$(env CC)", "gcc", "clang", "g++", "clang++")
     end
 
     return toolchains
@@ -174,10 +174,10 @@ function main(platform, name)
         -- check cuda
         check_cuda(config)
 
-        -- check cu-cxx after checking arch
+        -- check cu-ccbin after checking arch
         if config.get("cuda") then
             local toolchains = singleton.get("linux.toolchains", _toolchains)
-            check_toolchain(config, "cu-cxx", toolchains["cu-cxx"])
+            check_toolchain(config, "cu-ccbin", toolchains["cu-ccbin"])
         end
     end
 end
diff --git a/xmake/platforms/linux/global.lua b/xmake/platforms/linux/global.lua
index 4c8318a23..22aa9239a 100644
--- a/xmake/platforms/linux/global.lua
+++ b/xmake/platforms/linux/global.lua
@@ -20,7 +20,6 @@
 
 -- imports
 import("core.base.global")
-import("private.platform.check_cuda")
 
 -- check it
 function main(platform, name)
@@ -29,8 +28,5 @@ function main(platform, name)
     if name then
         raise("we cannot check global." .. name)
     end
-
-    -- check cuda
-    check_cuda(global)
 end
 
diff --git a/xmake/platforms/linux/load.lua b/xmake/platforms/linux/load.lua
index b9ff210b1..dcdacd8c1 100644
--- a/xmake/platforms/linux/load.lua
+++ b/xmake/platforms/linux/load.lua
@@ -77,22 +77,5 @@ function main(platform)
     -- init flags for rust
     platform:set("rc-shflags", "")
     platform:set("rc-ldflags", "")
-
-    -- init flags for cuda
-    local cu_archs = { i386 = "-m32 -Xcompiler -m32", x86_64 = "-m64 -Xcompiler -m64" }
-    platform:add("cuflags", cu_archs[arch] or "")
-    platform:add("cu-shflags", cu_archs[arch] or "")
-    platform:add("cu-ldflags", cu_archs[arch] or "")
-    local cuda_dir = config.get("cuda")
-    if cuda_dir then
-        platform:add("cuflags", "-I" .. os.args(path.join(cuda_dir, "include")))
-        platform:add("cu-ldflags", "-L" .. os.args(path.join(cuda_dir, "lib")))
-        platform:add("cu-shflags", "-L" .. os.args(path.join(cuda_dir, "lib")))
-        platform:add("cu-ldflags", "-Xlinker -rpath=" .. os.args(path.join(cuda_dir, "lib")))
-    end
-    local cu_cxx = config.get("cu-cxx")
-    if cu_cxx then
-        platform:add("cuflags", "-ccbin", os.args(cu_cxx))
-    end
 end
 
diff --git a/xmake/platforms/macosx/config.lua b/xmake/platforms/macosx/config.lua
index 59ad2d7bb..d2eaa807d 100644
--- a/xmake/platforms/macosx/config.lua
+++ b/xmake/platforms/macosx/config.lua
@@ -60,13 +60,13 @@ function _toolchains()
     local cu         = toolchain("the cuda compiler")
     local cu_ld      = toolchain("the cuda linker")
     local cu_sh      = toolchain("the cuda shared library linker")
-    local cu_cxx     = toolchain("the cuda host c++ compiler")
+    local cu_ccbin     = toolchain("the cuda host c++ compiler")
     local toolchains = {cc = cc, cxx = cxx, as = as, ld = ld, sh = sh, ar = ar, ex = ex, 
                         mm = mm, mxx = mxx, sc = sc, ["sc-ld"] = sc_ld, ["sc-sh"] = sc_sh,
                         gc = gc, ["gc-ld"] = gc_ld, ["gc-ar"] = gc_ar,
                         dc = dc, ["dc-ld"] = dc_ld, ["dc-sh"] = dc_sh, ["dc-ar"] = dc_ar,
                         rc = rc, ["rc-ld"] = rc_ld, ["rc-sh"] = rc_sh, ["rc-ar"] = rc_ar,
-                        cu = cu, ["cu-ld"] = cu_ld, ["cu-sh"] = cu_sh, ["cu-cxx"] = cu_cxx}
+                        cu = cu, ["cu-ld"] = cu_ld, ["cu-sh"] = cu_sh, ["cu-ccbin"] = cu_ccbin}
 
     -- init the c compiler
     cc:add("$(env CC)", {name = "clang", cross = cross}, "clang", "gcc")
@@ -133,7 +133,7 @@ function _toolchains()
     cu:add("nvcc")
     cu_ld:add("nvcc")
     cu_sh:add("nvcc")
-    cu_cxx:add("$(env CXX)", "$(env CC)", "clang", "gcc")
+    cu_ccbin:add("$(env CXX)", "$(env CC)", "clang", "gcc")
 
     return toolchains
 end
@@ -158,10 +158,10 @@ function main(platform, name)
         -- check cuda
         check_cuda(config)
 
-        -- check cu-cxx after checking arch
+        -- check cu-ccbin after checking arch
         if config.get("cuda") then
             local toolchains = singleton.get("macosx.toolchains." .. (config.get("arch") or os.arch()), _toolchains)
-            check_toolchain(config, "cu-cxx", toolchains["cu-cxx"])
+            check_toolchain(config, "cu-ccbin", toolchains["cu-ccbin"])
         end
     end
 end
diff --git a/xmake/platforms/macosx/global.lua b/xmake/platforms/macosx/global.lua
index 5b8ac731b..1db907fa7 100644
--- a/xmake/platforms/macosx/global.lua
+++ b/xmake/platforms/macosx/global.lua
@@ -20,7 +20,6 @@
 
 -- imports
 import("core.base.global")
-import("private.platform.check_cuda")
 import("private.platform.check_xcode")
 
 -- check it
@@ -33,8 +32,5 @@ function main(platform, name)
 
     -- check xcode 
     check_xcode(global, true)
-
-    -- check cuda
-    check_cuda(global)
 end
 
diff --git a/xmake/platforms/macosx/load.lua b/xmake/platforms/macosx/load.lua
index f47e7a31d..e53dc1b8c 100644
--- a/xmake/platforms/macosx/load.lua
+++ b/xmake/platforms/macosx/load.lua
@@ -88,23 +88,5 @@ function main(platform)
     -- init flags for rust
     platform:set("rc-shflags", "")
     platform:set("rc-ldflags", "")
-
-    -- init flags for cuda
-    local cuflags_arch = { i386 = "-m32 -Xcompiler -arch -Xcompiler i386", x86_64 = "-m64 -Xcompiler -arch -Xcompiler x86_64" }
-    local ldflags_arch = { i386 = "-m32 -Xlinker -arch -Xlinker i386", x86_64 = "-m64 -Xlinker -arch -Xlinker x86_64" }
-    platform:add("cuflags", cuflags_arch[arch] or "")
-    platform:add("cu-shflags", ldflags_arch[arch] or "")
-    platform:add("cu-ldflags", ldflags_arch[arch] or "")
-    local cuda_dir = config.get("cuda")
-    if cuda_dir then
-        platform:add("cuflags", "-I" .. os.args(path.join(cuda_dir, "include")))
-        platform:add("cu-ldflags", "-L" .. os.args(path.join(cuda_dir, "lib")))
-        platform:add("cu-shflags", "-L" .. os.args(path.join(cuda_dir, "lib")))
-        platform:add("cu-ldflags", "-Xlinker -rpath -Xlinker " .. os.args(path.join(cuda_dir, "lib")))
-    end
-    local cu_cxx = config.get("cu-cxx")
-    if cu_cxx then
-        platform:add("cuflags", "-ccbin", os.args(cu_cxx))
-    end
 end
 
diff --git a/xmake/platforms/windows/global.lua b/xmake/platforms/windows/global.lua
index 3068b81b7..2c5d7afda 100644
--- a/xmake/platforms/windows/global.lua
+++ b/xmake/platforms/windows/global.lua
@@ -21,7 +21,6 @@
 -- imports
 import("core.base.global")
 import("private.platform.check_arch")
-import("private.platform.check_cuda")
 import("private.platform.check_vstudio")
 
 -- clean temporary global configs
@@ -44,9 +43,6 @@ function main(platform, name)
     -- check vstudio
     check_vstudio(global)
 
-    -- check cuda
-    check_cuda(global)
-
     -- clean temporary global configs
     _clean_global()
 end
diff --git a/xmake/platforms/windows/load.lua b/xmake/platforms/windows/load.lua
index 2de1a138e..f7789ddb9 100644
--- a/xmake/platforms/windows/load.lua
+++ b/xmake/platforms/windows/load.lua
@@ -35,16 +35,4 @@ function main(platform)
     platform:add("dcflags", dc_archs[arch])
     platform:add("dc-shflags", dc_archs[arch])
     platform:add("dc-ldflags", dc_archs[arch])
-
-    -- init flags for cuda
-    local cu_archs = { x86 = "-m32", x64 = "-m64" }
-    platform:add("cuflags", cu_archs[arch] or "")
-    platform:add("cu-shflags", cu_archs[arch] or "")
-    platform:add("cu-ldflags", cu_archs[arch] or "")
-    local cuda_dir = config.get("cuda")
-    if cuda_dir then
-        platform:add("cuflags", "-I" .. os.args(path.join(cuda_dir, "include")))
-        platform:add("cu-ldflags", "-L" .. os.args(path.join(cuda_dir, "lib")))
-        platform:add("cu-shflags", "-L" .. os.args(path.join(cuda_dir, "lib")))
-    end
 end
diff --git a/xmake/rules/cuda/device_link/xmake.lua b/xmake/rules/cuda/device_link/xmake.lua
new file mode 100644
index 000000000..358c42d93
--- /dev/null
+++ b/xmake/rules/cuda/device_link/xmake.lua
@@ -0,0 +1,116 @@
+--!A cross-platform build utility based on Lua
+--
+-- 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.
+-- 
+-- Copyright (C) 2015 - 2019, TBOOX Open Source Group.
+--
+-- @author      ruki
+-- @file        xmake.lua
+--
+
+-- define rule: device-link
+rule("cuda.device_link")
+
+    -- add rule: cuda environment
+    add_deps("cuda.env")
+
+    -- clean files
+    after_clean(function (target)
+        os.tryrm(target:objectfile(path.join(".cuda", "devlink", target:basename() .. "_gpucode.cu")))
+        os.tryrm(target:dependfile(targetfile))
+    end)
+
+    -- @see https://devblogs.nvidia.com/separate-compilation-linking-cuda-device-code/
+    before_link(function (target, opt)
+
+        -- only for binary/shared
+        local targetkind = target:targetkind()
+        if targetkind ~= "binary" and targetkind ~= "shared" then
+            return 
+        end
+
+        -- imports
+        import("core.base.option")
+        import("core.theme.theme")
+        import("core.project.config")
+        import("core.project.depend")
+        import("core.tool.linker")
+
+        -- load linker instance
+        local linkinst = linker.load("gpucode", "cu", {target = target})
+
+        -- init culdflags
+        local culdflags = {"-dlink"}
+        if targetkind == "shared" then
+            table.insert(culdflags, "-shared")
+        end
+
+        -- get link flags
+        local linkflags = linkinst:linkflags({target = target, configs = {force = {culdflags = culdflags}}})
+
+        -- get target file
+        local targetfile = target:objectfile(path.join(".cuda", "devlink", target:basename() .. "_gpucode.cu"))
+
+        -- get object files
+        local objectfiles = nil
+        for sourcekind, sourcebatch in pairs(target:sourcebatches()) do
+            if sourcekind == "cu" then
+                objectfiles = sourcebatch.objectfiles
+            end
+        end
+        if not objectfiles then
+            return
+        end
+
+        -- insert gpucode.o to the object files
+        table.insert(target:objectfiles(), targetfile)
+
+        -- load dependent info 
+        local dependfile = target:dependfile(targetfile)
+        local dependinfo = option.get("rebuild") and {} or (depend.load(dependfile) or {})
+
+        -- need build this target?
+        local depfiles = objectfiles
+        local depvalues = {linkinst:program(), linkflags}
+        if not depend.is_changed(dependinfo, {lastmtime = os.mtime(target:targetfile()), values = depvalues, files = depfiles}) then
+            return 
+        end
+
+        -- is verbose?
+        local verbose = option.get("verbose")
+
+        -- trace progress info
+        cprintf("${color.build.progress}" .. theme.get("text.build.progress_format") .. ":${clear} ", opt.progress)
+        if verbose then
+            cprint("${dim color.build.target}devlinking.$(mode) %s", path.filename(targetfile))
+        else
+            cprint("${color.build.target}devlinking.$(mode) %s", path.filename(targetfile))
+        end
+
+        -- trace verbose info
+        if verbose then
+            print(linkinst:linkcmd(objectfiles, targetfile, {linkflags = linkflags}))
+        end
+
+        -- flush io buffer to update progress info
+        io.flush()
+
+        -- link it
+        assert(linkinst:link(objectfiles, targetfile, {linkflags = linkflags}))
+
+        -- update files and values to the dependent file
+        dependinfo.files  = depfiles
+        dependinfo.values = depvalues
+        depend.save(dependinfo, dependfile)
+    end)
+
diff --git a/xmake/rules/cuda/env/xmake.lua b/xmake/rules/cuda/env/xmake.lua
new file mode 100644
index 000000000..1acbc08ba
--- /dev/null
+++ b/xmake/rules/cuda/env/xmake.lua
@@ -0,0 +1,66 @@
+--!A cross-platform build utility based on Lua
+--
+-- 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.
+-- 
+-- Copyright (C) 2015 - 2019, TBOOX Open Source Group.
+--
+-- @author      ruki
+-- @file        xmake.lua
+--
+
+-- define rule: environment
+rule("cuda.env")
+    after_load(function (target)
+
+        -- get cuda sdk
+        import("detect.sdks.find_cuda")
+        local cuda = assert(find_cuda(nil, {verbose = true}), "Cuda SDK not found!")
+
+        -- add arch
+        if is_arch("i386", "x86") then
+            target:add("cuflags", "-m32", {force = true})
+            target:add("culdflags", "-m32", {force = true})
+        else
+            target:add("cuflags", "-m64", {force = true})
+            target:add("culdflags", "-m64", {force = true})
+        end
+
+        -- add -ccbin
+        local cu_ccbin = get_config("cu-ccbin")
+        if cu_ccbin then
+            target:add("cuflags", "-ccbin", os.args(cu_ccbin), {force = true})
+            target:add("culdflags", "-ccbin", os.args(cu_ccbin), {force = true})
+        end
+
+        -- add links
+        target:add("syslinks", "cudadevrt")
+        local cudart = false
+        for _, link in ipairs(table.join(target:get("links") or {}, target:get("syslinks"))) do
+            if link == "cudart" or link == "cudart_static" then
+                cudart = true
+                break
+            end
+        end
+        if not cudart then
+            target:add("syslinks", "cudart_static")
+        end
+        if is_plat("linux") then
+            target:add("syslinks", "rt", "pthread", "dl")
+        end
+        target:add("linkdirs", cuda.linkdirs)
+        target:add("rpathdirs", cuda.linkdirs)
+
+        -- add includedirs
+        target:add("includedirs", cuda.includedirs)
+    end)
+
diff --git a/xmake/includes/add_cugencodes.lua b/xmake/rules/cuda/gencodes/xmake.lua
index d68517962..9fcdbcd3a 100644
--- a/xmake/includes/add_cugencodes.lua
+++ b/xmake/rules/cuda/gencodes/xmake.lua
@@ -11,37 +11,29 @@
 -- 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.
---
+-- 
 -- Copyright (C) 2015 - 2019, TBOOX Open Source Group.
 --
--- @author      OpportunityLiu
--- @file        add_cugencodes.lua
---
-
--- add cuda `-gencode` flags to target
---
--- the gpu arch format syntax
--- - compute_xx                   --> `-gencode arch=compute_xx,code=compute_xx`
--- - sm_xx                        --> `-gencode arch=compute_xx,code=sm_xx`
--- - sm_xx,sm_yy                  --> `-gencode arch=compute_xx,code=[sm_xx,sm_yy]`
--- - compute_xx,sm_yy             --> `-gencode arch=compute_xx,code=sm_yy`
--- - compute_xx,sm_yy,sm_zz       --> `-gencode arch=compute_xx,code=[sm_yy,sm_zz]`
--- - native                       --> match the fastest cuda device on current host,
---                                    eg. for a Tesla P100, `-gencode arch=compute_60,code=sm_60` will be added,
---                                    if no available device is found, no `-gencode` flags will be added
---                                    @seealso xmake/modules/lib/detect/find_cudadevices
---
--- e.g.
--- includes("add_cugencodes.lua")
--- target("test")
---    set_kind("binary")
---    add_files("src/*.cu")
---    add_cugencodes("native", "compute_50,sm_50", "compute_70")
+-- @author      ruki
+-- @file        xmake.lua
 --
 
+-- define rule: gencodes
+rule("cuda.gencodes")
 
--- define rule
-rule("cuda.add_cugencodes")
+    -- add cuda `-gencode` flags to target
+    --
+    -- the gpu arch format syntax
+    -- - compute_xx                   --> `-gencode arch=compute_xx,code=compute_xx`
+    -- - sm_xx                        --> `-gencode arch=compute_xx,code=sm_xx`
+    -- - sm_xx,sm_yy                  --> `-gencode arch=compute_xx,code=[sm_xx,sm_yy]`
+    -- - compute_xx,sm_yy             --> `-gencode arch=compute_xx,code=sm_yy`
+    -- - compute_xx,sm_yy,sm_zz       --> `-gencode arch=compute_xx,code=[sm_yy,sm_zz]`
+    -- - native                       --> match the fastest cuda device on current host,
+    --                                    eg. for a Tesla P100, `-gencode arch=compute_60,code=sm_60` will be added,
+    --                                    if no available device is found, no `-gencode` flags will be added
+    --                                    @seealso xmake/modules/lib/detect/find_cudadevices
+    --
     before_load(function (target)
 
         local function set (list)
@@ -125,19 +117,16 @@ rule("cuda.add_cugencodes")
             end
         end
 
-        for _, v in ipairs(target:values("cuda.gencode")) do
+        local cugencodes = table.wrap(target:get("cugencodes"))
+        for _, opt in ipairs(target:orderopts()) do
+            table.join2(cugencodes, opt:get("cugencodes"))
+        end
+        for _, v in ipairs(cugencodes) do
             local flag = nf_cugencode(v)
             if flag then
                 target:add('cuflags', flag)
-                target:add('ldflags', flag)
+                target:add('culdflags', flag)
             end
         end
     end)
 rule_end()
-
--- add cuda gencode to target
-function add_cugencodes(...)
-    add_rules("cuda.add_cugencodes")
-    add_values("cuda.gencode", ...)
-end
-
diff --git a/xmake/rules/cuda/xmake.lua b/xmake/rules/cuda/xmake.lua
new file mode 100644
index 000000000..0ec1ddaee
--- /dev/null
+++ b/xmake/rules/cuda/xmake.lua
@@ -0,0 +1,26 @@
+--!A cross-platform build utility based on Lua
+--
+-- 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.
+-- 
+-- Copyright (C) 2015 - 2019, TBOOX Open Source Group.
+--
+-- @author      ruki
+-- @file        xmake.lua
+--
+
+-- define rule: cuda
+rule("cuda")
+
+    -- add rules
+    add_deps("cuda.device_link", "cuda.gencodes")
+
diff --git a/xmake/scripts/find_cudadevices.cu b/xmake/scripts/find_cudadevices.cpp
index 56d64ff2b..fd002cc84 100644
--- a/xmake/scripts/find_cudadevices.cu
+++ b/xmake/scripts/find_cudadevices.cpp
@@ -19,10 +19,19 @@ inline void check(cudaError_t result)
     }
 }
 
-inline void print_value(size_t value)
+inline void print_value(unsigned long long value)
 {
-    // in case we don't have '%zu'
-    printf("%llu", (unsigned long long)value);
+    printf("%llu", value);
+}
+
+inline void print_value(unsigned long value)
+{
+    printf("%lu", value);
+}
+
+inline void print_value(unsigned int value)
+{
+    printf("%u", value);
 }
 
 inline void print_value(bool value)
@@ -48,11 +57,6 @@ inline void print_value(const T (&value)[len])
     printf(")");
 }
 
-inline void print_value(unsigned int value)
-{
-    printf("%u", value);
-}
-
 inline void print_value(const void *value)
 {
     printf("\"%s\"", (const char *)value);
diff --git a/xmake/scripts/gitignore b/xmake/scripts/gitignore
index 08ffbbd04..152105761 100644
--- a/xmake/scripts/gitignore
+++ b/xmake/scripts/gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
+# Xmake cache
+.xmake/
+build/
+
+# MacOS Cache
 .DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+
+
diff --git a/xmake/templates/c++/console/project/.gitignore b/xmake/templates/c++/console/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c++/console/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c++/console_qt/project/.gitignore b/xmake/templates/c++/console_qt/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c++/console_qt/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c++/console_tbox/project/.gitignore b/xmake/templates/c++/console_tbox/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c++/console_tbox/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c++/shared_library/project/.gitignore b/xmake/templates/c++/shared_library/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c++/shared_library/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c++/shared_library_qt/project/.gitignore b/xmake/templates/c++/shared_library_qt/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c++/shared_library_qt/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c++/shared_library_tbox/project/.gitignore b/xmake/templates/c++/shared_library_tbox/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c++/shared_library_tbox/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c++/static_library/project/.gitignore b/xmake/templates/c++/static_library/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c++/static_library/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c++/static_library_qt/project/.gitignore b/xmake/templates/c++/static_library_qt/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c++/static_library_qt/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c++/static_library_tbox/project/.gitignore b/xmake/templates/c++/static_library_tbox/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c++/static_library_tbox/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c/console/project/.gitignore b/xmake/templates/c/console/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c/console/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c/console_tbox/.gitignore b/xmake/templates/c/console_tbox/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c/console_tbox/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c/shared_library/project/.gitignore b/xmake/templates/c/shared_library/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c/shared_library/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c/shared_library_tbox/project/.gitignore b/xmake/templates/c/shared_library_tbox/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c/shared_library_tbox/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c/static_library/project/.gitignore b/xmake/templates/c/static_library/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c/static_library/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/c/static_library_tbox/project/.gitignore b/xmake/templates/c/static_library_tbox/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/c/static_library_tbox/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/cuda/console/project/.gitignore b/xmake/templates/cuda/console/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/cuda/console/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/cuda/console/project/xmake.lua b/xmake/templates/cuda/console/project/xmake.lua
index 467ae7296..226488448 100644
--- a/xmake/templates/cuda/console/project/xmake.lua
+++ b/xmake/templates/cuda/console/project/xmake.lua
@@ -2,15 +2,17 @@
 -- add modes: debug and release
 add_rules("mode.debug", "mode.release")
 
--- add helper function add_cugencodes
-includes('add_cugencodes.lua')
-
 -- define target
 target("[targetname]")
 
     -- set kind
     set_kind("binary")
 
+    -- generate relocatable device code for device linker of dependents
+    -- if no __device__ and __global__ functions will be called cross file,
+    -- this instruction could be omitted
+    -- add_cuflags("-rdc=true")
+
     -- add files
     add_files("src/*.cu")
 
@@ -24,4 +26,4 @@ target("[targetname]")
     -- add_cugencodes("sm_30", "sm_35", "sm_37", "sm_50", "sm_52", "sm_60", "sm_61", "sm_70", "sm_75")
 
     -- -- generate PTX code from the highest SM architecture to guarantee forward-compatibility
-    -- add_cugencodes("compute_75")
\ No newline at end of file
+    -- add_cugencodes("compute_75")
diff --git a/xmake/templates/cuda/shared_library/project/inc/lib.cuh b/xmake/templates/cuda/shared_library/project/inc/lib.cuh
new file mode 100644
index 000000000..5565e3eb5
--- /dev/null
+++ b/xmake/templates/cuda/shared_library/project/inc/lib.cuh
@@ -0,0 +1,22 @@
+#pragma once
+
+#include "cuda_runtime.h"
+
+#ifdef __cplusplus
+extern "C"
+{
+#endif
+
+#if defined(_WIN32)
+#define __export __declspec(dllexport)
+#elif defined(__GNUC__) && ((__GNUC__ >= 4) || (__GNUC__ == 3 && __GNUC_MINOR__ >= 3))
+#define __export __attribute__((visibility("default")))
+#else
+#define __export
+#endif
+
+    __export cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);
+
+#ifdef __cplusplus
+}
+#endif
diff --git a/xmake/templates/cuda/shared_library/project/src/lib.cu b/xmake/templates/cuda/shared_library/project/src/lib.cu
new file mode 100644
index 000000000..d2cff9ac2
--- /dev/null
+++ b/xmake/templates/cuda/shared_library/project/src/lib.cu
@@ -0,0 +1,98 @@
+#include <lib.cuh>
+#include <stdio.h>
+#include "device_launch_parameters.h"
+
+__global__ void addKernel(int *c, const int *a, const int *b)
+{
+    int i = threadIdx.x;
+    c[i] = a[i] + b[i];
+}
+
+// Helper function for using CUDA to add vectors in parallel.
+cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
+{
+    int *dev_a = 0;
+    int *dev_b = 0;
+    int *dev_c = 0;
+    cudaError_t cudaStatus;
+
+    // Choose which GPU to run on, change this on a multi-GPU system.
+    cudaStatus = cudaSetDevice(0);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
+        goto Error;
+    }
+
+    // Allocate GPU buffers for three vectors (two input, one output)    .
+    cudaStatus = cudaMalloc((void **)&dev_c, size * sizeof(int));
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMalloc failed!");
+        goto Error;
+    }
+
+    cudaStatus = cudaMalloc((void **)&dev_a, size * sizeof(int));
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMalloc failed!");
+        goto Error;
+    }
+
+    cudaStatus = cudaMalloc((void **)&dev_b, size * sizeof(int));
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMalloc failed!");
+        goto Error;
+    }
+
+    // Copy input vectors from host memory to GPU buffers.
+    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMemcpy failed!");
+        goto Error;
+    }
+
+    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMemcpy failed!");
+        goto Error;
+    }
+
+    // Launch a kernel on the GPU with one thread for each element.
+    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
+
+    // Check for any errors launching the kernel
+    cudaStatus = cudaGetLastError();
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
+        goto Error;
+    }
+
+    // cudaDeviceSynchronize waits for the kernel to finish, and returns
+    // any errors encountered during the launch.
+    cudaStatus = cudaDeviceSynchronize();
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
+        goto Error;
+    }
+
+    // Copy output vector from GPU buffer to host memory.
+    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
+    if (cudaStatus != cudaSuccess)
+    {
+        fprintf(stderr, "cudaMemcpy failed!");
+        goto Error;
+    }
+
+Error:
+    cudaFree(dev_c);
+    cudaFree(dev_a);
+    cudaFree(dev_b);
+
+    return cudaStatus;
+}
diff --git a/xmake/templates/cuda/shared_library/project/xmake.lua b/xmake/templates/cuda/shared_library/project/xmake.lua
new file mode 100644
index 000000000..8ffb8b7e9
--- /dev/null
+++ b/xmake/templates/cuda/shared_library/project/xmake.lua
@@ -0,0 +1,35 @@
+
+-- add modes: debug and release
+add_rules("mode.debug", "mode.release")
+
+-- define target
+target("[targetname]")
+
+    -- set kind
+    set_kind("shared")
+
+    -- add modes: debug and release
+    add_rules("mode.debug", "mode.release")
+
+    -- add include dirs
+    add_includedirs("inc")
+
+    -- generate relocatable device code for device linker of dependents
+    -- if no __device__ and __global__ functions will be called cross file,
+    -- this instruction could be omitted
+    -- add_cuflags("-rdc=true")
+
+    -- add files
+    add_files("src/**.cu")
+
+    -- generate SASS code for SM architecture of current host
+    add_cugencodes("native")
+
+    -- generate PTX code for the virtual architecture to guarantee compatibility
+    add_cugencodes("compute_30")
+
+    -- -- generate SASS code for each SM architecture
+    -- add_cugencodes("sm_30", "sm_35", "sm_37", "sm_50", "sm_52", "sm_60", "sm_61", "sm_70", "sm_75")
+
+    -- -- generate PTX code from the highest SM architecture to guarantee forward-compatibility
+    -- add_cugencodes("compute_75")
diff --git a/xmake/templates/cuda/shared_library/template.lua b/xmake/templates/cuda/shared_library/template.lua
new file mode 100644
index 000000000..6df920968
--- /dev/null
+++ b/xmake/templates/cuda/shared_library/template.lua
@@ -0,0 +1,14 @@
+-- set name
+set_name("shared")
+
+-- set description
+set_description("The Shared Library")
+
+-- set project directory
+set_projectdir("project")
+
+-- add macros
+add_macros("targetname", "$(targetname)")
+
+-- add macro files
+add_macrofiles("xmake.lua")
\ No newline at end of file
diff --git a/xmake/templates/cuda/static_library/project/inc/lib.cuh b/xmake/templates/cuda/static_library/project/inc/lib.cuh
new file mode 100644
index 000000000..35255e31e
--- /dev/null
+++ b/xmake/templates/cuda/static_library/project/inc/lib.cuh
@@ -0,0 +1,3 @@
+#pragma once
+
+__global__ void addKernel(int *c, const int *a, const int *b);
\ No newline at end of file
diff --git a/xmake/templates/cuda/static_library/project/src/lib.cu b/xmake/templates/cuda/static_library/project/src/lib.cu
new file mode 100644
index 000000000..a5f054354
--- /dev/null
+++ b/xmake/templates/cuda/static_library/project/src/lib.cu
@@ -0,0 +1,7 @@
+#include <lib.cuh>
+
+__global__ void addKernel(int *c, const int *a, const int *b)
+{
+    int i = threadIdx.x;
+    c[i] = a[i] + b[i];
+}
\ No newline at end of file
diff --git a/xmake/templates/cuda/static_library/project/xmake.lua b/xmake/templates/cuda/static_library/project/xmake.lua
new file mode 100644
index 000000000..be08da0da
--- /dev/null
+++ b/xmake/templates/cuda/static_library/project/xmake.lua
@@ -0,0 +1,35 @@
+
+-- add modes: debug and release
+add_rules("mode.debug", "mode.release")
+
+-- define target
+target("[targetname]")
+
+    -- set kind
+    set_kind("static")
+
+    -- add modes: debug and release
+    add_rules("mode.debug", "mode.release")
+
+    -- add include dirs
+    add_includedirs("inc")
+
+    -- generate relocatable device code for device linker of dependents
+    -- if no __device__ and __global__ functions will be called cross file or be exported,
+    -- this instruction could be omitted
+    add_cuflags("-rdc=true")
+
+    -- add files
+    add_files("src/**.cu")
+
+    -- generate SASS code for SM architecture of current host
+    add_cugencodes("native")
+
+    -- generate PTX code for the virtual architecture to guarantee compatibility
+    add_cugencodes("compute_30")
+
+    -- -- generate SASS code for each SM architecture
+    -- add_cugencodes("sm_30", "sm_35", "sm_37", "sm_50", "sm_52", "sm_60", "sm_61", "sm_70", "sm_75")
+
+    -- -- generate PTX code from the highest SM architecture to guarantee forward-compatibility
+    -- add_cugencodes("compute_75")
diff --git a/xmake/templates/cuda/static_library/template.lua b/xmake/templates/cuda/static_library/template.lua
new file mode 100644
index 000000000..ff9c56179
--- /dev/null
+++ b/xmake/templates/cuda/static_library/template.lua
@@ -0,0 +1,14 @@
+-- set name
+set_name("static")
+
+-- set description
+set_description("The Static Library")
+
+-- set project directory
+set_projectdir("project")
+
+-- add macros
+add_macros("targetname", "$(targetname)")
+
+-- add macro files
+add_macrofiles("xmake.lua")
\ No newline at end of file
diff --git a/xmake/templates/dlang/console/project/.gitignore b/xmake/templates/dlang/console/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/dlang/console/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/dlang/shared_library/project/.gitignore b/xmake/templates/dlang/shared_library/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/dlang/shared_library/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/dlang/static_library/project/.gitignore b/xmake/templates/dlang/static_library/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/dlang/static_library/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/go/console/project/.gitignore b/xmake/templates/go/console/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/go/console/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/go/static_library/project/.gitignore b/xmake/templates/go/static_library/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/go/static_library/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/objc++/console/project/.gitignore b/xmake/templates/objc++/console/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/objc++/console/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/objc/console/project/.gitignore b/xmake/templates/objc/console/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/objc/console/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/rust/console/project/.gitignore b/xmake/templates/rust/console/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/rust/console/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/rust/static_library/project/.gitignore b/xmake/templates/rust/static_library/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/rust/static_library/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
diff --git a/xmake/templates/swift/console/project/.gitignore b/xmake/templates/swift/console/project/.gitignore
deleted file mode 100644
index 08ffbbd04..000000000
--- a/xmake/templates/swift/console/project/.gitignore
+++ /dev/null
@@ -1,22 +0,0 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build