diff options
Diffstat (limited to 'Tests/Cuda/Complex/file3.cu')
-rw-r--r-- | Tests/Cuda/Complex/file3.cu | 36 |
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; } |