cuda: add test for device linking

4 files changed, 73 insertions, 0 deletions

diff --git a/test cases/cuda/17 separate compilation linking/b.cu b/test cases/cuda/17 separate compilation linking/b.cu
new file mode 100644
index 0000000..33ff561
--- /dev/null
+++ b/test cases/cuda/17 separate compilation linking/b.cu
@@ -0,0 +1,5 @@
+#include "b.h"
+
+__device__ int g[N];
+
+__device__ void bar(void) { g[threadIdx.x]++; }
diff --git a/test cases/cuda/17 separate compilation linking/b.h b/test cases/cuda/17 separate compilation linking/b.h
new file mode 100644
index 0000000..d8a0efc
--- /dev/null
+++ b/test cases/cuda/17 separate compilation linking/b.h
@@ -0,0 +1,5 @@
+#define N 8
+
+extern __device__ int g[N];
+
+extern __device__ void bar(void);
diff --git a/test cases/cuda/17 separate compilation linking/main.cu b/test cases/cuda/17 separate compilation linking/main.cu
new file mode 100644
index 0000000..b07d01b
--- /dev/null
+++ b/test cases/cuda/17 separate compilation linking/main.cu
@@ -0,0 +1,44 @@
+#include <stdio.h>
+
+#include "b.h"
+
+__global__ void foo(void)
+{
+    __shared__ int a[N];
+    a[threadIdx.x] = threadIdx.x;
+
+    __syncthreads();
+
+    g[threadIdx.x] = a[blockDim.x - threadIdx.x - 1];
+
+    bar();
+}
+
+int main(void)
+{
+    unsigned int i;
+    int *dg, hg[N];
+    int sum = 0;
+
+    foo<<<1, N>>>();
+
+    if (cudaGetSymbolAddress((void**)&dg, g)) {
+        printf("couldn't get the symbol addr\n");
+        return 1;
+    }
+    if (cudaMemcpy(hg, dg, N * sizeof(int), cudaMemcpyDeviceToHost)) {
+        printf("couldn't memcpy\n");
+        return 1;
+    }
+
+    for (i = 0; i < N; i++) {
+        sum += hg[i];
+    }
+    if (sum == 36) {
+        printf("PASSED\n");
+    } else {
+        printf("FAILED (%d)\n", sum);
+    }
+
+    return 0;
+}
diff --git a/test cases/cuda/17 separate compilation linking/meson.build b/test cases/cuda/17 separate compilation linking/meson.build
new file mode 100644
index 0000000..ee86123
--- /dev/null
+++ b/test cases/cuda/17 separate compilation linking/meson.build
@@ -0,0 +1,19 @@
+# example here is inspired by Nvidia's blog post:
+#   https://developer.nvidia.com/blog/separate-compilation-linking-cuda-device-code/
+# code:
+#   https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#examples
+
+project('device linking', ['cpp', 'cuda'], version : '1.0.0')
+
+nvcc = meson.get_compiler('cuda')
+cuda = import('unstable-cuda')
+
+arch_flags = cuda.nvcc_arch_flags(nvcc.version(), 'Auto', detected : ['8.0'])
+
+message('NVCC version:   ' + nvcc.version())
+message('NVCC flags:     ' + ' '.join(arch_flags))
+
+# test device linking with -dc (which is equivalent to `--relocatable-device-code true`)
+lib = static_library('devicefuncs', ['b.cu'], cuda_args : ['-dc'] + arch_flags)
+exe = executable('app', 'main.cu', cuda_args : ['-dc'] + arch_flags, link_with : lib, link_args : arch_flags)
+test('cudatest', exe)