summaryrefslogtreecommitdiffstats
path: root/Tests/Cuda/Complex/file3.cu
diff options
context:
space:
mode:
Diffstat (limited to 'Tests/Cuda/Complex/file3.cu')
-rw-r--r--Tests/Cuda/Complex/file3.cu36
1 files changed, 29 insertions, 7 deletions
diff --git a/Tests/Cuda/Complex/file3.cu b/Tests/Cuda/Complex/file3.cu
index 7c37d66..d055b42 100644
--- a/Tests/Cuda/Complex/file3.cu
+++ b/Tests/Cuda/Complex/file3.cu
@@ -7,20 +7,42 @@
result_type __device__ file1_func(int x);
result_type_dynamic __device__ file2_func(int x);
-static __global__ void file3_kernel(result_type& r, int x)
+static __global__ void file3_kernel(result_type* r, int x)
{
- r = file1_func(x);
+ *r = file1_func(x);
result_type_dynamic rd = file2_func(x);
}
int file3_launch_kernel(int x)
{
- result_type r;
+ result_type* r;
+ cudaError_t err = cudaMallocManaged(&r, sizeof(result_type));
+ if (err != cudaSuccess) {
+ std::cerr << "file3_launch_kernel: cudaMallocManaged failed: "
+ << cudaGetErrorString(err) << std::endl;
+ return x;
+ }
+
file3_kernel<<<1, 1>>>(r, x);
- cudaError_t err = cudaGetLastError();
- if (err == cudaSuccess) {
- std::cerr << cudaGetErrorString(err) << std::endl;
+ err = cudaGetLastError();
+ if (err != cudaSuccess) {
+ std::cerr << "file3_kernel [SYNC] failed: " << cudaGetErrorString(err)
+ << std::endl;
+ return x;
+ }
+ err = cudaDeviceSynchronize();
+ if (err != cudaSuccess) {
+ std::cerr << "file3_kernel [ASYNC] failed: "
+ << cudaGetErrorString(cudaGetLastError()) << std::endl;
+ return x;
+ }
+ int result = r->sum;
+ err = cudaFree(r);
+ if (err != cudaSuccess) {
+ std::cerr << "file3_launch_kernel: cudaFree failed: "
+ << cudaGetErrorString(err) << std::endl;
return x;
}
- return r.sum;
+
+ return result;
}