Merge pull request #450 from OpportunityLiu/cuda-tests

Cuda tests & templates

47 files changed, 813 insertions, 635 deletions

diff --git a/.gitignore b/.gitignore
index cd7ecbe72..99bc3124f 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,44 +1,26 @@
-*.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
+
+!/xmake/actions/build/
\ No newline at end of file
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/scripts/gitignore b/xmake/scripts/gitignore
index 08ffbbd04..b849bb8f7 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
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c++/console/project/.gitignore b/xmake/templates/c++/console/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c++/console/project/.gitignore
+++ b/xmake/templates/c++/console/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c++/console_qt/project/.gitignore b/xmake/templates/c++/console_qt/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c++/console_qt/project/.gitignore
+++ b/xmake/templates/c++/console_qt/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c++/console_tbox/project/.gitignore b/xmake/templates/c++/console_tbox/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c++/console_tbox/project/.gitignore
+++ b/xmake/templates/c++/console_tbox/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c++/shared_library/project/.gitignore b/xmake/templates/c++/shared_library/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c++/shared_library/project/.gitignore
+++ b/xmake/templates/c++/shared_library/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c++/shared_library_qt/project/.gitignore b/xmake/templates/c++/shared_library_qt/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c++/shared_library_qt/project/.gitignore
+++ b/xmake/templates/c++/shared_library_qt/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c++/shared_library_tbox/project/.gitignore b/xmake/templates/c++/shared_library_tbox/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c++/shared_library_tbox/project/.gitignore
+++ b/xmake/templates/c++/shared_library_tbox/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c++/static_library/project/.gitignore b/xmake/templates/c++/static_library/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c++/static_library/project/.gitignore
+++ b/xmake/templates/c++/static_library/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c++/static_library_qt/project/.gitignore b/xmake/templates/c++/static_library_qt/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c++/static_library_qt/project/.gitignore
+++ b/xmake/templates/c++/static_library_qt/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c++/static_library_tbox/project/.gitignore b/xmake/templates/c++/static_library_tbox/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c++/static_library_tbox/project/.gitignore
+++ b/xmake/templates/c++/static_library_tbox/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c/console/project/.gitignore b/xmake/templates/c/console/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c/console/project/.gitignore
+++ b/xmake/templates/c/console/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c/console_tbox/.gitignore b/xmake/templates/c/console_tbox/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c/console_tbox/.gitignore
+++ b/xmake/templates/c/console_tbox/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c/shared_library/project/.gitignore b/xmake/templates/c/shared_library/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c/shared_library/project/.gitignore
+++ b/xmake/templates/c/shared_library/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c/shared_library_tbox/project/.gitignore b/xmake/templates/c/shared_library_tbox/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c/shared_library_tbox/project/.gitignore
+++ b/xmake/templates/c/shared_library_tbox/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c/static_library/project/.gitignore b/xmake/templates/c/static_library/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c/static_library/project/.gitignore
+++ b/xmake/templates/c/static_library/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/c/static_library_tbox/project/.gitignore b/xmake/templates/c/static_library_tbox/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/c/static_library_tbox/project/.gitignore
+++ b/xmake/templates/c/static_library_tbox/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/cuda/console/project/.gitignore b/xmake/templates/cuda/console/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/cuda/console/project/.gitignore
+++ b/xmake/templates/cuda/console/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/cuda/console/project/xmake.lua b/xmake/templates/cuda/console/project/xmake.lua
index 1c1c48cfb..226488448 100644
--- a/xmake/templates/cuda/console/project/xmake.lua
+++ b/xmake/templates/cuda/console/project/xmake.lua
@@ -8,6 +8,11 @@ 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")
 
diff --git a/xmake/templates/cuda/shared_library/project/.gitignore b/xmake/templates/cuda/shared_library/project/.gitignore
new file mode 100644
index 000000000..173fb6eaf
--- /dev/null
+++ b/xmake/templates/cuda/shared_library/project/.gitignore
@@ -0,0 +1,2 @@
+.xmake
+build
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/.gitignore b/xmake/templates/cuda/static_library/project/.gitignore
new file mode 100644
index 000000000..173fb6eaf
--- /dev/null
+++ b/xmake/templates/cuda/static_library/project/.gitignore
@@ -0,0 +1,2 @@
+.xmake
+build
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
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/dlang/console/project/.gitignore
+++ b/xmake/templates/dlang/console/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/dlang/shared_library/project/.gitignore b/xmake/templates/dlang/shared_library/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/dlang/shared_library/project/.gitignore
+++ b/xmake/templates/dlang/shared_library/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/dlang/static_library/project/.gitignore b/xmake/templates/dlang/static_library/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/dlang/static_library/project/.gitignore
+++ b/xmake/templates/dlang/static_library/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/go/console/project/.gitignore b/xmake/templates/go/console/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/go/console/project/.gitignore
+++ b/xmake/templates/go/console/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/go/static_library/project/.gitignore b/xmake/templates/go/static_library/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/go/static_library/project/.gitignore
+++ b/xmake/templates/go/static_library/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/objc++/console/project/.gitignore b/xmake/templates/objc++/console/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/objc++/console/project/.gitignore
+++ b/xmake/templates/objc++/console/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/objc/console/project/.gitignore b/xmake/templates/objc/console/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/objc/console/project/.gitignore
+++ b/xmake/templates/objc/console/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/rust/console/project/.gitignore b/xmake/templates/rust/console/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/rust/console/project/.gitignore
+++ b/xmake/templates/rust/console/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/rust/static_library/project/.gitignore b/xmake/templates/rust/static_library/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/rust/static_library/project/.gitignore
+++ b/xmake/templates/rust/static_library/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+
diff --git a/xmake/templates/swift/console/project/.gitignore b/xmake/templates/swift/console/project/.gitignore
index 08ffbbd04..b849bb8f7 100644
--- a/xmake/templates/swift/console/project/.gitignore
+++ b/xmake/templates/swift/console/project/.gitignore
@@ -1,22 +1,8 @@
-*.a
-*.o
-*.exe
-*.obj
-*.dll
-*.lib
-*.out
-*.suo
-*~
-*.swp
-*.swo
-*.bak
-*.orig
-*.pdb
-*.idb
-.svn
-.DS_Store
-.xmake
-*.gch
-*.gch.d
-gmon.out
-build
+# xmake temp files
+.xmake/
+# build result files
+build/
+
+.DS_Store/
+
+