Merge pull request #455 from OpportunityLiu/clang-cuda

Support clang as cuda compiler

36 files changed, 494 insertions, 139 deletions

diff --git a/CHANGELOG.md b/CHANGELOG.md
index 46ad3e30b..5311ccee8 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -4,6 +4,7 @@
 
 ### New features
 
+* [#455](https://github.com/xmake-io/xmake/pull/455): support clang as cuda compiler, try `xmake f --cu=clang`
 * [#440](https://github.com/xmake-io/xmake/issues/440): Add `set_rundir()` and `add_runenvs()` api for target/run
 * [#443](https://github.com/xmake-io/xmake/pull/443): Add tab completion support
 * Add `on_link`, `before_link` and `after_link` for rule and target
@@ -12,7 +13,7 @@
 ### Changes
 
 * [#430](https://github.com/xmake-io/xmake/pull/430): Add `add_cucodegens()` api to improve set codegen for cuda
-* [#432](https://github.com/xmake-io/xmake/pull/432): support deps analyze for cu file 
+* [#432](https://github.com/xmake-io/xmake/pull/432): support deps analyze for cu file (for CUDA 10.1+)
 * [#437](https://github.com/xmake-io/xmake/issues/437): Support explict git source for xmake update, `xmake update github:xmake-io/xmake#dev`
 * [#438](https://github.com/xmake-io/xmake/pull/438): Support to only update scripts, `xmake update --scriptonly dev`
 * [#433](https://github.com/xmake-io/xmake/issues/433): Improve cuda to support device-link
@@ -602,6 +603,7 @@
 
 ### 新特性
 
+* [#455](https://github.com/xmake-io/xmake/pull/455): 支持使用 clang 作为 cuda 编译器，`xmake f --cu=clang`
 * [#440](https://github.com/xmake-io/xmake/issues/440): 为target/run添加`set_rundir()`和`add_runenvs()`接口设置
 * [#443](https://github.com/xmake-io/xmake/pull/443): 添加命令行tab自动完成支持
 * 为rule/target添加`on_link`,`before_link`和`after_link`阶段自定义脚本支持
@@ -610,7 +612,7 @@
 ### 改进
 
 * [#430](https://github.com/xmake-io/xmake/pull/430): 添加`add_cucodegens()`api为cuda改进设置codegen
-* [#432](https://github.com/xmake-io/xmake/pull/432): 针对cuda编译支持依赖分析检测
+* [#432](https://github.com/xmake-io/xmake/pull/432): 针对cuda编译支持依赖分析检测（仅支持 CUDA 10.1+）
 * [#437](https://github.com/xmake-io/xmake/issues/437): 支持指定更新源，`xmake update github:xmake-io/xmake#dev`
 * [#438](https://github.com/xmake-io/xmake/pull/438): 支持仅更新脚本，`xmake update --scriptonly dev`
 * [#433](https://github.com/xmake-io/xmake/issues/433): 改进cuda构建支持device-link设备代码链接
diff --git a/tests/modules/path/test.lua b/tests/modules/path/test.lua
new file mode 100644
index 000000000..4217dec16
--- /dev/null
+++ b/tests/modules/path/test.lua
@@ -0,0 +1,27 @@
+function test_splitenv_win(t)
+    if not is_host("windows") then
+        return t:skip("wrong host platform")
+    end
+    t:are_equal(path.splitenv(""), {})
+    t:are_equal(path.splitenv("a"), {'a'})
+    t:are_equal(path.splitenv("a;b"), {'a','b'})
+    t:are_equal(path.splitenv(";;a;;b;"), {'a','b'})
+    t:are_equal(path.splitenv('c:/a;c:\\b'), {'c:/a', 'c:\\b'})
+    t:are_equal(path.splitenv('"a;aa;aa;;"'), {"a;aa;aa;;"})
+    t:are_equal(path.splitenv('"a;aa;aa;;";bb;;'), {"a;aa;aa;;", 'bb'})
+    t:are_equal(path.splitenv('"a;aa;aa;;";"a;cc;aa;;";bb;"d";'), {"a;aa;aa;;","a;cc;aa;;", 'bb', 'd' })
+end
+
+function test_splitenv_unix(t)
+    if is_host("windows") then
+        return t:skip("wrong host platform")
+    end
+    t:are_equal(path.splitenv(""), {})
+    t:are_equal(path.splitenv("a"), {'a'})
+    t:are_equal(path.splitenv("a:b"), {'a','b'})
+    t:are_equal(path.splitenv("::a::b:"), {'a','b'})
+    t:are_equal(path.splitenv('a%tag:b'), {'a','b'})
+    t:are_equal(path.splitenv('a%tag:b%tag'), {'a','b'})
+    t:are_equal(path.splitenv('a%tag:b%%tag%%'), {'a','b'})
+    t:are_equal(path.splitenv('a%tag:b:%tag:'), {'a','b'})
+end
\ No newline at end of file
diff --git a/tests/projects/cuda/console_2/inc/lib.cuh b/tests/projects/cuda/console_2/inc/lib.cuh
new file mode 100644
index 000000000..35255e31e
--- /dev/null
+++ b/tests/projects/cuda/console_2/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/console_2/src/lib.cu b/tests/projects/cuda/console_2/src/lib.cu
new file mode 100644
index 000000000..a5f054354
--- /dev/null
+++ b/tests/projects/cuda/console_2/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/console_2/src/main.cu b/tests/projects/cuda/console_2/src/main.cu
new file mode 100644
index 000000000..32844cd56
--- /dev/null
+++ b/tests/projects/cuda/console_2/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/console_2/xmake.lua b/tests/projects/cuda/console_2/xmake.lua
new file mode 100644
index 000000000..2c445d7e3
--- /dev/null
+++ b/tests/projects/cuda/console_2/xmake.lua
@@ -0,0 +1,19 @@
+
+-- 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("bin")
+
+    -- set kind
+    set_kind("binary")
+
+    add_cuflags("-rdc=true")
+
+    add_includedirs("inc")
+
+    -- add files
+    add_files("src/*.cu")
diff --git a/tests/projects/dlang/console/test.lua b/tests/projects/dlang/console/test.lua
index b423189da..c1bac30ee 100644
--- a/tests/projects/dlang/console/test.lua
+++ b/tests/projects/dlang/console/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build()
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/projects/dlang/shared_library/test.lua b/tests/projects/dlang/shared_library/test.lua
index b423189da..c1bac30ee 100644
--- a/tests/projects/dlang/shared_library/test.lua
+++ b/tests/projects/dlang/shared_library/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build()
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/projects/dlang/static_library/test.lua b/tests/projects/dlang/static_library/test.lua
index b423189da..c1bac30ee 100644
--- a/tests/projects/dlang/static_library/test.lua
+++ b/tests/projects/dlang/static_library/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build()
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/projects/go/console/test.lua b/tests/projects/go/console/test.lua
index b423189da..c1bac30ee 100644
--- a/tests/projects/go/console/test.lua
+++ b/tests/projects/go/console/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build()
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/projects/go/static_library/test.lua b/tests/projects/go/static_library/test.lua
index b423189da..c1bac30ee 100644
--- a/tests/projects/go/static_library/test.lua
+++ b/tests/projects/go/static_library/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build()
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/projects/objc++/console/test.lua b/tests/projects/objc++/console/test.lua
index f73079781..60277b691 100644
--- a/tests/projects/objc++/console/test.lua
+++ b/tests/projects/objc++/console/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build({iphoneos = true})
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/projects/objc/console/test.lua b/tests/projects/objc/console/test.lua
index f73079781..60277b691 100644
--- a/tests/projects/objc/console/test.lua
+++ b/tests/projects/objc/console/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build({iphoneos = true})
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/projects/rust/console/test.lua b/tests/projects/rust/console/test.lua
index b423189da..c1bac30ee 100644
--- a/tests/projects/rust/console/test.lua
+++ b/tests/projects/rust/console/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build()
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/projects/rust/static_library/test.lua b/tests/projects/rust/static_library/test.lua
index b423189da..c1bac30ee 100644
--- a/tests/projects/rust/static_library/test.lua
+++ b/tests/projects/rust/static_library/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build()
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/projects/swift/console/test.lua b/tests/projects/swift/console/test.lua
index f73079781..60277b691 100644
--- a/tests/projects/swift/console/test.lua
+++ b/tests/projects/swift/console/test.lua
@@ -4,5 +4,7 @@ function main(t)
     -- build project
     if os.host() == "macosx" then
         t:build({iphoneos = true})
+    else
+        return t:skip("wrong host platform")
     end
 end
diff --git a/tests/runner.lua b/tests/runner.lua
index b4325edc2..227e66a7f 100644
--- a/tests/runner.lua
+++ b/tests/runner.lua
@@ -1,7 +1,5 @@
 import("core.base.option")
-import("test_utils.test_assert")
-import("test_utils.print_error")
-import("test_utils.build", { alias = "test_build" })
+import("test_utils.context", { alias = "test_context" })
 
 function main(script)
 
@@ -16,14 +14,7 @@ function main(script)
     os.setenv("XMAKE_STATS", "false")
 
     -- init test context
-    local context =
-    {
-        filename = script
-    }
-    context = test_assert(context)
-    function context:build(argv)
-        test_build(argv)
-    end
+    local context = test_context(script)
 
     local root = path.directory(script)
 
@@ -50,10 +41,10 @@ function main(script)
             if verbose then print(">>     running %s ...", k) end
             context.func = v
             context.funcname = k
-            try
+            local result = try
             {
                 function ()
-                    v(context)
+                    return v(context)
                 end,
                 catch
                 {
@@ -66,6 +57,9 @@ function main(script)
                     end
                 }
             }
+            if context:is_skipped(result) then
+                print(">>     skipped %s : %s", k, result.reason)
+            end
             succeed_count = succeed_count + 1
         end
     end
diff --git a/tests/test_utils/context.lua b/tests/test_utils/context.lua
new file mode 100644
index 000000000..50271ac76
--- /dev/null
+++ b/tests/test_utils/context.lua
@@ -0,0 +1,13 @@
+import("test_build")
+import("test_skip")
+import("test_assert")
+
+function main(filename)
+
+    local context = { filename = filename }
+    table.join2(context, test_build())
+    table.join2(context, test_skip())
+    table.join2(context, test_assert())
+
+    return context
+end
\ No newline at end of file
diff --git a/tests/test_utils/test_assert.lua b/tests/test_utils/test_assert.lua
index 4a97cb3ac..a964fcfb3 100644
--- a/tests/test_utils/test_assert.lua
+++ b/tests/test_utils/test_assert.lua
@@ -1,8 +1,8 @@
 
-import("print_error")
+
 import("check")
 
-local test_assert = { print_error = print_error.main }
+local test_assert = { print_error = import("print_error", { anonymous = true }).main }
 
 function test_assert:require(value)
     if not value then
@@ -66,7 +66,6 @@ function test_assert:will_raise(func, message_pattern)
     table.remove(self._will_raise_stack, 1)
 end
 
-function main(context)
-    table.join2(context, test_assert)
-    return context
+function main()
+    return test_assert
 end
diff --git a/tests/test_utils/build.lua b/tests/test_utils/test_build.lua
index 7725f1ddc..4625a5943 100644
--- a/tests/test_utils/build.lua
+++ b/tests/test_utils/test_build.lua
@@ -1,8 +1,9 @@
 -- imports
 import("privilege.sudo")
 
--- main entry
-function main(argv)
+local test_build = {}
+
+function test_build:build(argv)
 
     -- check global config
     os.exec("xmake g -c")
@@ -35,3 +36,7 @@ function main(argv)
         end
     end
 end
+
+function main()
+    return test_build
+end
diff --git a/tests/test_utils/test_skip.lua b/tests/test_utils/test_skip.lua
new file mode 100644
index 000000000..c17f2c42c
--- /dev/null
+++ b/tests/test_utils/test_skip.lua
@@ -0,0 +1,13 @@
+local test_skip = { _is_skipped_tag = {true} }
+
+function test_skip:skip(reason)
+    return { is_skipped = self._is_skipped_tag, reason = reason, context = self }
+end
+
+function test_skip:is_skipped(result)
+    return result and result.context and result.context._is_skipped_tag == result.is_skipped
+end
+
+function main()
+    return test_skip
+end
\ No newline at end of file
diff --git a/xmake/core/base/path.lua b/xmake/core/base/path.lua
index 489f6b702..c5ba7423b 100644
--- a/xmake/core/base/path.lua
+++ b/xmake/core/base/path.lua
@@ -27,6 +27,10 @@ local string    = require("base/string")
 
 -- get the directory of the path
 function path.directory(p)
+
+    -- check
+    assert(p)
+
     local i = p:find_last("[/\\]")
     if i then
         if i > 1 then i = i - 1 end
@@ -38,6 +42,10 @@ end
 
 -- get the filename of the path
 function path.filename(p)
+
+    -- check
+    assert(p)
+
     local i = p:find_last("[/\\]")
     if i then
         return p:sub(i + 1)
@@ -48,6 +56,10 @@ end
 
 -- get the basename of the path
 function path.basename(p)
+
+    -- check
+    assert(p)
+
     local name = path.filename(p)
     local i = name:find_last(".", true)
     if i then
@@ -89,6 +101,10 @@ end
 
 -- split path by the separator
 function path.split(p)
+
+    -- check
+    assert(p)
+
     return p:split("/\\")
 end
 
@@ -102,15 +118,60 @@ function path.envsep()
     return xmake._HOST == "windows" and ';' or ':'
 end
 
+-- split environment variable with `path.envsep()`,
+-- also handles more speical cases such as posix flags and windows quoted pathes
+function path.splitenv(env_path)
+
+    -- check
+    assert(env_path)
+
+    local result = {}
+    if xmake._HOST == "windows" then
+        while #env_path > 0 do
+            if env_path:startswith(path.envsep()) then
+                env_path = env_path:sub(2)
+            elseif env_path:startswith('"') then
+                -- path quoted with, can contain `;`
+                local p_end = env_path:find('"' .. path.envsep(), 2, true) or env_path:find('"$', 2) or (#env_path + 1)
+                table.insert(result, env_path:sub(2, p_end - 1))
+                env_path = env_path:sub(p_end + 1)
+            else
+                local p_end = env_path:find(path.envsep(), 2, true) or (#env_path + 1)
+                table.insert(result, env_path:sub(1, p_end - 1))
+                env_path = env_path:sub(p_end)
+            end
+        end
+    else
+        -- see https://git.kernel.org/pub/scm/utils/dash/dash.git/tree/src/exec.c?h=v0.5.9.1&id=afe0e0152e4dc12d84be3c02d6d62b0456d68580#n173
+        -- no escape sequences, so `:` and `%` is invalid in environment variable
+        for _, v in ipairs(env_path:split(path.envsep(), { plain = true })) do
+            -- flag for shells, style `<path>%<flag>`
+            local flag = v:find("%", 1, true)
+            if flag then
+                v = v:sub(1, flag - 1)
+            end
+            if #v > 0 then
+                table.insert(result, v)
+            end
+        end
+    end
+
+    return result
+end
+
 -- the last character is the path seperator?
 function path.islastsep(p)
+
+    -- check
+    assert(p)
+
     local sep = p:sub(#p, #p)
     return xmake._HOST == "windows" and (sep == '\\' or sep == '/') or (sep == '/')
 end
 
 -- convert path pattern to a lua pattern
 function path.pattern(pattern)
-    
+
     -- translate wildcards, .e.g *, **
     pattern = pattern:gsub("([%+%.%-%^%$%(%)%%])", "%%%1")
     pattern = pattern:gsub("%*%*", "\001")
diff --git a/xmake/core/sandbox/modules/import/lib/detect/find_file.lua b/xmake/core/sandbox/modules/import/lib/detect/find_file.lua
index 1543d8bd6..2d64d6067 100644
--- a/xmake/core/sandbox/modules/import/lib/detect/find_file.lua
+++ b/xmake/core/sandbox/modules/import/lib/detect/find_file.lua
@@ -75,31 +75,37 @@ function sandbox_lib_detect_find_file.main(name, pathes, opt)
 
         -- format path for builtin variables
         if type(_path) == "function" then
-            local ok, results = sandbox.load(_path) 
+            local ok, results = sandbox.load(_path)
             if ok then
                 _path = results or ""
-            else 
+            else
                 raise(results)
             end
-        else
-            _path = vformat(_path)
+        elseif type(_path) == "string" then
+            if _path:match("^%$%($s*env%s+%S+%s*%)$") then
+                _path = path.splitenv(vformat(_path))
+            else
+                _path = vformat(_path)
+            end
         end
 
-        -- find file with suffixes
-        if #suffixes > 0 then
-            for _, suffix in ipairs(suffixes) do
-                local filedir = path.join(_path, suffix)
-                local results = sandbox_lib_detect_find_file._find(filedir, name)
+        for _, _s_path in ipairs(table.wrap(_path)) do
+            -- find file with suffixes
+            if #suffixes > 0 then
+                for _, suffix in ipairs(suffixes) do
+                    local filedir = path.join(_s_path, suffix)
+                    local results = sandbox_lib_detect_find_file._find(filedir, name)
+                    if results then
+                        return results
+                    end
+                end
+            else
+                -- find file in the given path
+                local results = sandbox_lib_detect_find_file._find(_s_path, name)
                 if results then
                     return results
                 end
             end
-        else
-            -- find file in the given path
-            local results = sandbox_lib_detect_find_file._find(_path, name)
-            if results then
-                return results
-            end
         end
     end
 end
diff --git a/xmake/core/sandbox/modules/import/lib/detect/find_program.lua b/xmake/core/sandbox/modules/import/lib/detect/find_program.lua
index d9cfbf484..bfc21f35c 100644
--- a/xmake/core/sandbox/modules/import/lib/detect/find_program.lua
+++ b/xmake/core/sandbox/modules/import/lib/detect/find_program.lua
@@ -63,7 +63,7 @@ function sandbox_lib_detect_find_program._check(program, opt)
     if type(opt.check) == "string" then
         ok, errors = os.runv(program, {opt.check})
     else
-        ok, errors = sandbox.load(opt.check, program) 
+        ok, errors = sandbox.load(opt.check, program)
     end
 
     -- check failed? print verbose error info
@@ -84,29 +84,36 @@ function sandbox_lib_detect_find_program._find_from_pathes(name, pathes, opt)
 
             -- format path for builtin variables
             if type(_path) == "function" then
-                local ok, results = sandbox.load(_path) 
+                local ok, results = sandbox.load(_path)
                 if ok then
                     _path = results or ""
-                else 
+                else
                     raise(results)
                 end
-            else
-                _path = vformat(_path)
+            elseif type(_path) == "string" then
+                if _path:match("^%$%($s*env%s+%S+%s*%)$") then
+                    _path = path.splitenv(vformat(_path))
+                else
+                    _path = vformat(_path)
+                end
             end
 
-            -- get program path
-            local program_path = nil
-            if os.isfile(_path) then
-                program_path = _path
-            elseif os.isdir(_path) then
-                program_path = path.join(_path, name)
-            end
+            for _, _s_path in ipairs(table.wrap(_path)) do
 
-            -- the program path
-            if program_path and (os.isexec(program_path) or os.isexec(program_path:split("%s")[1])) then
-                -- check it
-                if sandbox_lib_detect_find_program._check(program_path, opt) then
-                    return program_path
+                -- get program path
+                local program_path = nil
+                if os.isfile(_s_path) then
+                    program_path = _s_path
+                elseif os.isdir(_s_path) then
+                    program_path = path.join(_s_path, name)
+                end
+
+                -- the program path
+                if program_path and (os.isexec(program_path) or os.isexec(program_path:split("%s")[1])) then
+                    -- check it
+                    if sandbox_lib_detect_find_program._check(program_path, opt) then
+                        return program_path
+                    end
                 end
             end
         end
@@ -119,7 +126,7 @@ function sandbox_lib_detect_find_program._find_from_packages(name, opt)
     -- get the manifest file of package, .e.g ~/.xmake/packages/g/git/1.1.12/ed41d5327fad3fc06fe376b4a94f62ef/manifest.txt 
     local manifest_file = path.join(package.installdir(), name:sub(1, 1), name, opt.version, opt.buildhash, "manifest.txt")
     if not os.isfile(manifest_file) then
-        return 
+        return
     end
 
     -- get install directory of this package
@@ -250,7 +257,7 @@ function sandbox_lib_detect_find_program.main(name, opt)
     end
 
     -- attempt to get result from cache first
-    local cacheinfo = cache.load(cachekey) 
+    local cacheinfo = cache.load(cachekey)
     local result = cacheinfo[name]
     if result ~= nil and not opt.force then
         return utils.ifelse(result, result, nil)
@@ -258,7 +265,7 @@ function sandbox_lib_detect_find_program.main(name, opt)
 
     -- find executable program
     checking = utils.ifelse(coroutine_running, name, nil)
-    result = sandbox_lib_detect_find_program._find(name, opt.pathes, opt) 
+    result = sandbox_lib_detect_find_program._find(name, opt.pathes, opt)
     checking = nil
 
     -- cache result
diff --git a/xmake/languages/cuda/xmake.lua b/xmake/languages/cuda/xmake.lua
index 75b4bf5b8..0f532d4ed 100644
--- a/xmake/languages/cuda/xmake.lua
+++ b/xmake/languages/cuda/xmake.lua
@@ -140,14 +140,11 @@ language("cuda")
                     {category = "Cross Complation Configuration/Compiler Configuration"         }
                 ,   {nil, "cu",         "kv", nil,          "The Cuda 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"                   }
-                ,   {nil, "cu-ar",      "kv", nil,          "The Cuda Static Library Archiver"  }
-                ,   {nil, "cu-sh",      "kv", nil,          "The Cuda Shared Library Linker"    }
 
                 ,   {category = "Cross Complation Configuration/Compiler Flags Configuration"   }
                 ,   {nil, "cuflags",    "kv", nil,          "The Cuda Compiler Flags"           }
+                ,   {nil, "culdflags",  "kv", nil,          "The Cuda Linker Flags"             }
 
                 ,   {category = "Cross Complation Configuration/Builtin Flags Configuration"    }
                 ,   {nil, "links",      "kv", nil,          "The Link Libraries"                }
diff --git a/xmake/modules/core/tools/clang.lua b/xmake/modules/core/tools/clang.lua
index b7291b285..a57dd924f 100644
--- a/xmake/modules/core/tools/clang.lua
+++ b/xmake/modules/core/tools/clang.lua
@@ -23,15 +23,26 @@ inherit("gcc")
 
 -- init it
 function init(self)
-    
+
     -- init super
     _super.init(self)
 
+    if not is_plat("windows", "mingw") then
+        self:add("shared.cuflags", "-fPIC")
+    end
+
     -- suppress warning 
     self:add("cxflags", "-Qunused-arguments")
+    self:add("cuflags", "-Qunused-arguments")
     self:add("mxflags", "-Qunused-arguments")
     self:add("asflags", "-Qunused-arguments")
 
+    local cuda = get_config("cuda")
+    if cuda then
+        local cuda_path = "--cuda-path=" .. os.args(path.translate(cuda))
+        self:add("cuflags", cuda_path)
+    end
+
     -- init flags map
     self:set("mapflags",
     {
@@ -44,6 +55,16 @@ function init(self)
          -- strip
     ,   ["-s"]  = "-s"
     ,   ["-S"]  = "-S"
+
+        -- rdc
+    ,   ["-rdc=true"] = "-fcuda-rdc"
+    ,   ["-rdc true"] = "-fcuda-rdc"
+    ,   ["--relocatable-device-code=true"] = "-fcuda-rdc"
+    ,   ["--relocatable-device-code true"] = "-fcuda-rdc"
+    ,   ["-rdc=false"] = ""
+    ,   ["-rdc false"] = ""
+    ,   ["--relocatable-device-code=false"] = ""
+    ,   ["--relocatable-device-code false"] = ""
     })
 
 end
diff --git a/xmake/modules/core/tools/nvcc.lua b/xmake/modules/core/tools/nvcc.lua
index 592634ac8..e72857a93 100644
--- a/xmake/modules/core/tools/nvcc.lua
+++ b/xmake/modules/core/tools/nvcc.lua
@@ -29,11 +29,17 @@ import("private.tools.nvcc.parse_deps")
 -- init it
 function init(self)
 
-    -- init flags
-    if not is_plat("windows") then
+    -- init cuflags
+    if not is_plat("windows", "mingw") then
         self:set("shared.cuflags", "-Xcompiler -fPIC")
     end
 
+    -- add -ccbin
+    local cu_ccbin = get_config("cu-ccbin")
+    if cu_ccbin then
+        self:add("cuflags", "-ccbin=" .. os.args(cu_ccbin))
+    end
+
     -- init flags map
     self:set("mapflags",
     {
diff --git a/xmake/modules/detect/sdks/find_cuda.lua b/xmake/modules/detect/sdks/find_cuda.lua
index f9077a485..b7fb4911e 100644
--- a/xmake/modules/detect/sdks/find_cuda.lua
+++ b/xmake/modules/detect/sdks/find_cuda.lua
@@ -37,6 +37,7 @@ function _find_sdkdir()
     else
         table.insert(pathes, "/usr/local/cuda*/bin")
     end
+    table.insert(pathes, "$(env PATH)")
 
     -- attempt to find nvcc
     local nvcc = find_file(os.host() == "windows" and "nvcc.exe" or "nvcc", pathes)
diff --git a/xmake/modules/detect/tools/find_nvcc.lua b/xmake/modules/detect/tools/find_nvcc.lua
index 5b00510c1..5cd11e3dc 100644
--- a/xmake/modules/detect/tools/find_nvcc.lua
+++ b/xmake/modules/detect/tools/find_nvcc.lua
@@ -43,8 +43,12 @@ function main(opt)
     opt       = opt or {}
     opt.parse = opt.parse or "V(%d+%.?%d*%.?%d*.-)%s"
 
+    local program = nil
+
     -- find program
-    local program = find_program(opt.program or "nvcc", opt)
+    if opt.program then
+        program = find_program(opt.program, opt)
+    end
 
     -- not found? attempt to find program from cuda toolchains
     if not program then
@@ -54,9 +58,14 @@ function main(opt)
         end
     end
 
+    -- not found? attempt to find program from PATH
+    if not program then
+        program = find_program("nvcc", opt)
+    end
+
     -- find program version
     local version = nil
-    if program and opt and opt.version then
+    if program and opt.version then
         version = find_programver(program, opt)
     end
 
diff --git a/xmake/modules/lib/detect/find_cudadevices.lua b/xmake/modules/lib/detect/find_cudadevices.lua
index aa44fe0fb..63874685d 100644
--- a/xmake/modules/lib/detect/find_cudadevices.lua
+++ b/xmake/modules/lib/detect/find_cudadevices.lua
@@ -23,6 +23,8 @@ import("core.base.option")
 import("core.platform.platform")
 import("core.project.config")
 import("lib.detect.cache")
+import("lib.detect.find_tool")
+import("detect.sdks.find_cuda")
 
 -- a magic string to filter output
 local _PRINT_SUFFIX = "<find_cudadevices>"
@@ -30,7 +32,7 @@ local _PRINT_SUFFIX = "<find_cudadevices>"
 -- filter stdout and stderr with _PRINT_SUFFIX
 function _get_lines(str)
     local result = {}
-    for _, l in ipairs(str:split('\n')) do
+    for _, l in ipairs(str:split("\n")) do
         if l:startswith(_PRINT_SUFFIX) then
             table.insert(result, l:sub(#_PRINT_SUFFIX + 1))
         end
@@ -55,8 +57,8 @@ function _parse_value(value)
         return value:sub(2, -2)
     end
 
-    if value:startswith('(') and value:endswith(')') then
-        local values = value:sub(2, -2):split(',')
+    if value:startswith("(") and value:endswith(")") then
+        local values = value:sub(2, -2):split(",")
         local result = {}
         for _, v in ipairs(values) do
             table.insert(result, _parse_value(v:trim()))
@@ -77,7 +79,7 @@ function _parse_line(line, device)
     if key and value then
         key = key:trim()
         value = value:trim()
-        assert(not device[key], 'duplicate key: ' .. key)
+        assert(not device[key], "duplicate key: " .. key)
         device[key] = _parse_value(value)
     end
 end
@@ -99,7 +101,7 @@ function _parse_result(lines, verbose)
         end
         local devId = tonumber(l:match("%s*DEVICE #(%d+)"))
         if devId then
-            currentDevice = { ['$id'] = devId }
+            currentDevice = { ["$id"] = devId }
             table.insert(devices, currentDevice)
         elseif currentDevice then
             _parse_line(l, currentDevice)
@@ -111,32 +113,28 @@ end
 
 -- find devices
 function _find_devices(verbose)
-
-    local nvcc = platform.tool("cu")
-    if nvcc == nil then
-        raise('nvcc not found')
-    end
+    local nvcc = assert(find_tool("nvcc"), "nvcc not found")
 
     if verbose then
         cprint("${dim}checking for cuda devices")
     end
 
-    local sourcefile = path.join(os.programdir(), 'scripts', 'find_cudadevices.cpp')
+    local sourcefile = path.join(os.programdir(), "scripts", "find_cudadevices.cpp")
     local outfile = os.tmpfile()
+    local args = { sourcefile, "-run", "-o", outfile , '-DPRINT_SUFFIX="' .. _PRINT_SUFFIX .. '"' }
+
     local compile_errors = nil
-    local results, errors = try 
-    { 
+    local results, errors = try
+    {
         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 
+            return os.iorunv(nvcc.program, args)
+        end,
+        catch
         {
-            function (errs) 
-                compile_errors = tostring(errs) 
+            function (errs)
+                compile_errors = tostring(errs)
             end
-        } 
+        }
     }
 
     if compile_errors then
@@ -149,13 +147,13 @@ function _find_devices(verbose)
 
     -- clean up
     os.tryrm(outfile)
-    os.tryrm(outfile .. '.*')
+    os.tryrm(outfile .. ".*")
 
     -- get results
     local results_lines = _get_lines(results)
     local errors_lines = _get_lines(errors)
     if #errors_lines ~= 0 then
-        utils.warning("failed to find cuda devices: " .. table.concat(errors_lines, '\n'))
+        utils.warning("failed to find cuda devices: " .. table.concat(errors_lines, "\n"))
         return nil
     end
 
@@ -163,7 +161,7 @@ function _find_devices(verbose)
     local devices = _parse_result(results_lines, option.get("diagnosis"))
     if verbose then
         for _, v in ipairs(devices) do
-            cprint("${dim}> found device #%d: ${green bright}%s${reset dim} with compute ${bright}%d.%d${reset dim} capability", v['$id'], v.name, v.major, v.minor)
+            cprint("${dim}> found device #%d: ${green bright}%s${reset dim} with compute ${bright}%d.%d${reset dim} capability", v["$id"], v.name, v.major, v.minor)
         end
     end
     return devices
@@ -244,10 +242,10 @@ function _order_by_flops(devices)
         else
             sm_per_multiproc = ngpu_arch_cores_per_sm[dev.major * 10 + dev.minor] or 64;
         end
-        dev['$flops'] = dev.multiProcessorCount * sm_per_multiproc * dev.clockRate
+        dev["$flops"] = dev.multiProcessorCount * sm_per_multiproc * dev.clockRate
     end
 
-    table.sort(devices, function (a,b) return a['$flops'] > b['$flops'] end)
+    table.sort(devices, function (a,b) return a["$flops"] > b["$flops"] end)
     return devices
 end
 
@@ -256,7 +254,7 @@ end
 -- @param opt   the options
 --              e.g. { verbose = false, force = false, cachekey = "xxxx", min_sm_arch = 35, skip_compute_mode_prohibited = false, order_by_flops = true }
 --
--- @return      { { ['$id'] = 0, name = "GeForce GTX 960M", major = 5, minor = 0, ... }, ... }
+-- @return      { { ["$id"] = 0, name = "GeForce GTX 960M", major = 5, minor = 0, ... }, ... }
 --              for all keys, see https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaDeviceProp.html#structcudaDeviceProp
 --              keys might be differ as your cuda version varies
 --
diff --git a/xmake/platforms/linux/config.lua b/xmake/platforms/linux/config.lua
index ebccfc985..c86144804 100644
--- a/xmake/platforms/linux/config.lua
+++ b/xmake/platforms/linux/config.lua
@@ -77,14 +77,13 @@ function _toolchains()
     local rc_ar      = toolchain("the rust static library archiver")
     local cu         = toolchain("the cuda compiler")
     local cu_ld      = toolchain("the cuda linker")
-    local cu_sh      = toolchain("the cuda shared library linker")
     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-ccbin"] = cu_ccbin}
+                        cu = cu, ["cu-ld"] = cu_ld, ["cu-ccbin"] = cu_ccbin}
 
     -- init the c compiler
     cc:add("$(env CC)", {name = "gcc", cross = cross}, {name = "clang", cross = cross})
@@ -147,9 +146,8 @@ function _toolchains()
     rc_ar:add("$(env RC)", "rustc")
 
     -- init the cuda compiler and linker
-    cu:add("nvcc")
+    cu:add("nvcc", "clang++", "clang")
     cu_ld:add("nvcc")
-    cu_sh:add("nvcc")
     if not cross or cross == "" then
         cu_ccbin:add("$(env CXX)", "$(env CC)", "gcc", "clang", "g++", "clang++")
     end
diff --git a/xmake/platforms/macosx/config.lua b/xmake/platforms/macosx/config.lua
index a5b1935ba..3db232687 100644
--- a/xmake/platforms/macosx/config.lua
+++ b/xmake/platforms/macosx/config.lua
@@ -59,14 +59,13 @@ function _toolchains()
     local rc_ar      = toolchain("the rust static library archiver")
     local cu         = toolchain("the cuda compiler")
     local cu_ld      = toolchain("the cuda linker")
-    local cu_sh      = toolchain("the cuda shared library linker")
     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-ccbin"] = cu_ccbin}
+                        cu = cu, ["cu-ld"] = cu_ld, ["cu-ccbin"] = cu_ccbin}
 
     -- init the c compiler
     cc:add("$(env CC)", {name = "clang", cross = cross}, "clang", "gcc")
@@ -130,9 +129,8 @@ function _toolchains()
     rc_ar:add("$(env RC)", "rustc")
 
     -- init the cuda compiler and linker
-    cu:add("nvcc")
+    cu:add("nvcc", "clang")
     cu_ld:add("nvcc")
-    cu_sh:add("nvcc")
     cu_ccbin:add("$(env CXX)", "$(env CC)", "clang", "gcc")
 
     return toolchains
diff --git a/xmake/platforms/windows/config.lua b/xmake/platforms/windows/config.lua
index c61a1ab1a..1ba24cb65 100644
--- a/xmake/platforms/windows/config.lua
+++ b/xmake/platforms/windows/config.lua
@@ -56,12 +56,11 @@ function _toolchains()
     local rc_ar      = toolchain("the rust static library archiver")
     local cu         = toolchain("the cuda compiler")
     local cu_ld      = toolchain("the cuda linker")
-    local cu_sh      = toolchain("the cuda shared library linker")
     local toolchains = {cc = cc, cxx = cxx, mrc = mrc, as = as, ld = ld, sh = sh, ar = ar, ex = ex, 
                         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 = cu, ["cu-ld"] = cu_ld}
 
     -- init the c compiler
     cc:add("cl.exe")
@@ -109,9 +108,8 @@ function _toolchains()
     rc_ar:add("$(env RC)", "rustc")
 
     -- init the cuda compiler and linker
-    cu:add("nvcc")
+    cu:add("nvcc", "clang")
     cu_ld:add("nvcc")
-    cu_sh:add("nvcc")
 
     return toolchains
 end
diff --git a/xmake/rules/cuda/devlink/xmake.lua b/xmake/rules/cuda/devlink/xmake.lua
index 0416b9abe..1cdfb2fb7 100644
--- a/xmake/rules/cuda/devlink/xmake.lua
+++ b/xmake/rules/cuda/devlink/xmake.lua
@@ -26,16 +26,17 @@ rule("cuda.devlink")
 
     -- @see https://devblogs.nvidia.com/separate-compilation-linking-cuda-device-code/
     before_link(function (target, opt)
+        import("core.platform.platform")
 
         -- disable devlink?
         if target:values("cuda.devlink") == false then
-            return 
+            return
         end
 
         -- only for binary/shared
         local targetkind = target:targetkind()
         if targetkind ~= "binary" and targetkind ~= "shared" then
-            return 
+            return
         end
 
         -- imports
@@ -50,6 +51,8 @@ rule("cuda.devlink")
 
         -- init culdflags
         local culdflags = {"-dlink"}
+
+        -- add shared flag
         if targetkind == "shared" then
             table.insert(culdflags, "-shared")
         end
diff --git a/xmake/rules/cuda/env/xmake.lua b/xmake/rules/cuda/env/xmake.lua
index 1acbc08ba..a78e13af1 100644
--- a/xmake/rules/cuda/env/xmake.lua
+++ b/xmake/rules/cuda/env/xmake.lua
@@ -35,11 +35,10 @@ rule("cuda.env")
             target:add("culdflags", "-m64", {force = true})
         end
 
-        -- add -ccbin
+        -- 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})
+            target:add("culdflags", "-ccbin=" .. os.args(cu_ccbin), {force = true})
         end
 
         -- add links
diff --git a/xmake/rules/cuda/gencodes/xmake.lua b/xmake/rules/cuda/gencodes/xmake.lua
index 9fcdbcd3a..cecf83694 100644
--- a/xmake/rules/cuda/gencodes/xmake.lua
+++ b/xmake/rules/cuda/gencodes/xmake.lua
@@ -36,6 +36,9 @@ rule("cuda.gencodes")
     --
     before_load(function (target)
 
+        import("core.platform.platform")
+        import("lib.detect.find_cudadevices")
+
         local function set (list)
             local result = {}
             for _, l in ipairs(list) do result[l] = true end
@@ -43,17 +46,15 @@ rule("cuda.gencodes")
         end
 
         -- sm_20 and compute_20 is supported until CUDA 8
-        local knownVArchs = set { 20, 30, 32, 35, 37, 50, 52, 53, 60, 61, 62, 70, 72, 75, }
-        local knownRArchs = set { 20, 30, 32, 35, 37, 50, 52, 53, 60, 61, 62, 70, 72, 75, }
+        local known_v_archs = set { 20, 30, 32, 35, 37, 50, 52, 53, 60, 61, 62, 70, 72, 75, }
+        local known_r_archs = set { 20, 30, 32, 35, 37, 50, 52, 53, 60, 61, 62, 70, 72, 75, }
 
         local function nf_cugencode(archs)
-
             if type(archs) ~= 'string' then
                 return nil
             end
             archs = archs:trim():lower()
             if archs == 'native' then
-                import("lib.detect.find_cudadevices")
                 local device = find_cudadevices({ skip_compute_mode_prohibited = true, order_by_flops = true })[1]
                 if device then
                     return nf_cugencode('sm_' .. device.major .. device.minor)
@@ -61,10 +62,10 @@ rule("cuda.gencodes")
                 return nil
             end
 
-            local vArch = nil
-            local rArchs = {}
+            local v_arch = nil
+            local r_archs = {}
 
-            local function parse_arch(value, prefix, knowList)
+            local function parse_arch(value, prefix, know_list)
                 if not value:startswith(prefix) then
                     return nil
                 end
@@ -72,8 +73,8 @@ rule("cuda.gencodes")
                 if arch == nil then
                     raise("Unknown architecture: " .. value)
                 end
-                if not knowList[arch] then
-                    if arch <= table.maxn(knowList) then
+                if not know_list[arch] then
+                    if arch <= table.maxn(know_list) then
                         raise("Unknown architecture: " .. prefix .. "_" .. arch)
                     else
                         utils.warning("Unknown architecture: " .. prefix .. "_" .. arch)
@@ -84,37 +85,55 @@ rule("cuda.gencodes")
 
             for _, v in ipairs(archs:split(',')) do
                 local arch = v:trim()
-                local tempRArch = parse_arch(arch, 'sm', knownRArchs)
-                if tempRArch then
-                    table.insert(rArchs, tempRArch)
+                local temp_r_arch = parse_arch(arch, 'sm', known_r_archs)
+                if temp_r_arch then
+                    table.insert(r_archs, temp_r_arch)
                 end
 
-                local tempVArch = parse_arch(arch, 'compute', knownVArchs)
-                if tempVArch then
-                    if vArch ~= nil then
-                        raise("More than one virtual architecture is defined in one gpu gencode option: compute_" .. vArch .. " and compute_" .. tempVArch)
+                local temp_v_arch = parse_arch(arch, 'compute', known_v_archs)
+                if temp_v_arch then
+                    if v_arch ~= nil then
+                        raise("More than one virtual architecture is defined in one gpu gencode option: compute_" .. v_arch .. " and compute_" .. temp_v_arch)
                     end
-                    vArch = tempVArch
+                    v_arch = temp_v_arch
                 end
-                if not (tempRArch or tempVArch) then
+                if not (temp_r_arch or temp_v_arch) then
                     raise("Unknown architecture: " .. arch)
                 end
             end
 
-            if vArch == nil and #rArchs == 0 then
+            local result = { clang = {}, nvcc = {} }
+            if v_arch == nil and #r_archs == 0 then
                 return nil
             end
-            if #rArchs == 0 then
-                return '-gencode arch=compute_' .. vArch .. ',code=compute_' .. vArch
+
+            if #r_archs == 0 then
+                return {
+                    clang = '--cuda-gpu-arch=sm_' .. v_arch
+                ,   nvcc = '-gencode arch=compute_' .. v_arch .. ',code=compute_' .. v_arch }
+            end
+
+            if v_arch then
+                table.insert(r_archs, v_arch)
+            end
+            r_archs = table.unique(r_archs)
+            local clang_flags = {}
+            for _, r_arch in ipairs(r_archs) do
+                table.insert(clang_flags, '--cuda-gpu-arch=sm_' .. r_arch)
             end
 
-            rArchs = table.unique(rArchs)
-            vArch = vArch or math.min(unpack(rArchs))
-            if #rArchs == 1 then
-                return '-gencode arch=compute_' .. vArch .. ',code=sm_' .. rArchs[1]
+            r_archs = table.unique(r_archs)
+            v_arch = v_arch or math.min(unpack(r_archs))
+            local nvcc_flags = nil
+            if #r_archs == 1 then
+                nvcc_flags = '-gencode arch=compute_' .. v_arch .. ',code=sm_' .. r_archs[1]
             else
-                return '-gencode arch=compute_' .. vArch .. ',code=[sm_' .. table.concat(rArchs, ',sm_') .. ']'
+                nvcc_flags = '-gencode arch=compute_' .. v_arch .. ',code=[sm_' .. table.concat(r_archs, ',sm_') .. ']'
             end
+
+            return {
+                clang = clang_flags
+            ,   nvcc = nvcc_flags }
         end
 
         local cugencodes = table.wrap(target:get("cugencodes"))
@@ -124,8 +143,13 @@ rule("cuda.gencodes")
         for _, v in ipairs(cugencodes) do
             local flag = nf_cugencode(v)
             if flag then
-                target:add('cuflags', flag)
-                target:add('culdflags', flag)
+                local tool, toolname = platform.tool("cu")
+                if (toolname or path.basename(tool)) == "nvcc" then
+                    target:add('cuflags', flag.nvcc)
+                else
+                    target:add('cuflags', flag.clang)
+                end
+                target:add('culdflags', flag.nvcc)
             end
         end
     end)