shfl_scan 's shuffle_integral_image_test has some misleading code.

https://github.com/NVIDIA/cuda-samples/blob/5f97d7d0dff880bc6567faa4c5e62e389a6d6999/Samples/2_Concepts_and_Techniques/shfl_scan/shfl_scan.cu#L329-L333

The size of d_data and d_integral_image here can be misleading.

  int n_elements = w * h;
  int sz = sizeof(unsigned int) * n_elements;
  checkCudaErrors(cudaMallocHost(reinterpret_cast<void **>(&h_image), sz));
  
  
  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_data), sz));
  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_integral_image),
                             n_elements * sizeof(int) * 4));
  checkCudaErrors(cudaMemset(d_data, 1, sz));
  checkCudaErrors(cudaMemset(d_integral_image, 0, sz));

d_data 's size is (w * h) *sizeof(char) in this ap, while malloc it with w * h *sizeof(unsigned int).
d_integral_image's size is n_elements * sizeof(int) in this ap.
Current version has a puzzling size.

change it to:

  int n_elements = w * h;
  int sz = sizeof(unsigned int) * n_elements;
  checkCudaErrors(cudaMallocHost(reinterpret_cast<void **>(&h_image), sz));
  
  
  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_data), sz / 4));
  checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_integral_image),
                             n_elements * sizeof(int)));
  checkCudaErrors(cudaMemset(d_data, 1, sz / 4));
  checkCudaErrors(cudaMemset(d_integral_image, 0, sz));

may be better?

Open Samples/2_Concepts_and_Techniques/shfl_scan/shfl_scan.cu around lines 329-333 and inspect how d_data and d_integral_image are used after allocation. Compare their allocated and memset sizes with the element types and dimensions in the sample. Done means the allocation and initialization sizes clearly match the buffers' actual data.