Support clang as cuda compiler

15 files changed, 289 insertions, 52 deletions

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/xmake/core/base/path.lua b/xmake/core/base/path.lua
index 489f6b702..9499799f1 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
 
@@ -104,13 +120,17 @@ 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/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..1ddcdc562 100644
--- a/xmake/modules/core/tools/nvcc.lua
+++ b/xmake/modules/core/tools/nvcc.lua
@@ -29,11 +29,20 @@ import("private.tools.nvcc.parse_deps")
 -- init it
 function init(self)
 
-    -- init flags
+    -- init culdflags
+    self:set("shared.culdflags", "-shared")
+
+    -- init cuflags
     if not is_plat("windows") 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/lib/detect/find_cudadevices.lua b/xmake/modules/lib/detect/find_cudadevices.lua
index aa44fe0fb..8df34568e 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>"
@@ -112,7 +114,9 @@ end
 -- find devices
 function _find_devices(verbose)
 
-    local nvcc = platform.tool("cu")
+    local cuda = find_cuda(get_config("cuda"))
+    local nvcc = find_tool("nvcc", { program = path.join(cuda.bindir, "nvcc") })
+
     if nvcc == nil then
         raise('nvcc not found')
     end
@@ -124,19 +128,19 @@ function _find_devices(verbose)
     local sourcefile = path.join(os.programdir(), 'scripts', 'find_cudadevices.cpp')
     local outfile = os.tmpfile()
     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, { sourcefile, arch, '-run', '-o', outfile , '-DPRINT_SUFFIX="' .. _PRINT_SUFFIX .. '"' })
+        end,
+        catch
         {
-            function (errs) 
-                compile_errors = tostring(errs) 
+            function (errs)
+                compile_errors = tostring(errs)
             end
-        } 
+        }
     }
 
     if compile_errors then
diff --git a/xmake/platforms/linux/config.lua b/xmake/platforms/linux/config.lua
index ebccfc985..4f133120f 100644
--- a/xmake/platforms/linux/config.lua
+++ b/xmake/platforms/linux/config.lua
@@ -84,7 +84,7 @@ function _toolchains()
                         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})
diff --git a/xmake/platforms/macosx/config.lua b/xmake/platforms/macosx/config.lua
index a5b1935ba..2a8a4e080 100644
--- a/xmake/platforms/macosx/config.lua
+++ b/xmake/platforms/macosx/config.lua
@@ -66,7 +66,7 @@ function _toolchains()
                         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")
diff --git a/xmake/platforms/windows/config.lua b/xmake/platforms/windows/config.lua
index c61a1ab1a..fecc9d909 100644
--- a/xmake/platforms/windows/config.lua
+++ b/xmake/platforms/windows/config.lua
@@ -61,7 +61,7 @@ function _toolchains()
                         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")
diff --git a/xmake/rules/cuda/devlink/xmake.lua b/xmake/rules/cuda/devlink/xmake.lua
index 0416b9abe..2d4644781 100644
--- a/xmake/rules/cuda/devlink/xmake.lua
+++ b/xmake/rules/cuda/devlink/xmake.lua
@@ -26,16 +26,21 @@ 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?
+        -- local cu_tool, cu_toolname = platform.tool("cu")
+        -- if (cu_toolname or path.basename(cu_tool)) ~= "nvcc" then
+        --     return
+        -- end
         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 +55,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)