diff --git a/benchmark/benchmark.cu b/benchmark/benchmark.cu
index 51d22d7..77d4559 100644
--- a/benchmark/benchmark.cu
+++ b/benchmark/benchmark.cu
@@ -593,10 +593,12 @@ int main(int argc, char** argv) {
     // Note: excluding scaling from timing
 #ifdef R2C
     scale<<<(pinfo_x_r.size + 1024 - 1) / 1024, 1024>>>(output_r, 1.0 / fftsize, pinfo_x_r);
+    CHECK_CUDA_LAUNCH_EXIT();
     if (out_of_place) std::swap(input, output);
     if (out_of_place) std::swap(input_r, output_r);
 #else
     scale<<<(pinfo_x_c.size + 1024 - 1) / 1024, 1024>>>(output, 1.0 / fftsize, pinfo_x_c);
+    CHECK_CUDA_LAUNCH_EXIT();
     if (out_of_place) std::swap(input, output);
 #endif
 
diff --git a/examples/cc/basic_usage/basic_usage.cu b/examples/cc/basic_usage/basic_usage.cu
index 63249d9..d84fe8a 100644
--- a/examples/cc/basic_usage/basic_usage.cu
+++ b/examples/cc/basic_usage/basic_usage.cu
@@ -48,6 +48,15 @@
     }                                                                                                                  \
   } while (false)
 
+#define CHECK_CUDA_LAUNCH_EXIT()                                                                                       \
+  do {                                                                                                                 \
+    cudaError_t err = cudaGetLastError();                                                                              \
+    if (cudaSuccess != err) {                                                                                          \
+      fprintf(stderr, "%s:%d CUDA error. (%s)\n", __FILE__, __LINE__, cudaGetErrorString(err));                        \
+      exit(EXIT_FAILURE);                                                                                              \
+    }                                                                                                                  \
+  } while (false)
+
 #define CHECK_MPI_EXIT(call)                                                                                           \
   {                                                                                                                    \
     int err = call;                                                                                                    \
@@ -223,6 +232,7 @@ int main(int argc, char** argv) {
   int threads_per_block = 256;
   int nblocks = (pinfo_x.size + threads_per_block - 1) / threads_per_block;
   initialize_pencil<<<nblocks, threads_per_block>>>(data_d, pinfo_x);
+  CHECK_CUDA_LAUNCH_EXIT();
 
   // Allocating cuDecomp workspace
 
diff --git a/examples/cc/basic_usage/basic_usage_autotune.cu b/examples/cc/basic_usage/basic_usage_autotune.cu
index 3f76d14..1da7f75 100644
--- a/examples/cc/basic_usage/basic_usage_autotune.cu
+++ b/examples/cc/basic_usage/basic_usage_autotune.cu
@@ -48,6 +48,15 @@
     }                                                                                                                  \
   } while (false)
 
+#define CHECK_CUDA_LAUNCH_EXIT()                                                                                       \
+  do {                                                                                                                 \
+    cudaError_t err = cudaGetLastError();                                                                              \
+    if (cudaSuccess != err) {                                                                                          \
+      fprintf(stderr, "%s:%d CUDA error. (%s)\n", __FILE__, __LINE__, cudaGetErrorString(err));                        \
+      exit(EXIT_FAILURE);                                                                                              \
+    }                                                                                                                  \
+  } while (false)
+
 #define CHECK_MPI_EXIT(call)                                                                                           \
   {                                                                                                                    \
     int err = call;                                                                                                    \
@@ -207,6 +216,7 @@ int main(int argc, char** argv) {
   int threads_per_block = 256;
   int nblocks = (pinfo_x.size + threads_per_block - 1) / threads_per_block;
   initialize_pencil<<<nblocks, threads_per_block>>>(data_d, pinfo_x);
+  CHECK_CUDA_LAUNCH_EXIT();
 
   // Allocating cuDecomp workspace