Errors when executing code in pre-Workshop run through

[Sparrow0hawk]

I've been working through this workshop in prep for delivering the course and noticed a couple of issues that i'll document here and suggest PRs for.

1.

• [x] numbers isn't defined in https://carpentries-incubator.github.io/lesson-gpu-programming/03-numba/index.html for final code block

  %timeit -n 10 -r 1 check_prime_gpu(numbers)

  I think we mean't to define it in a cell above where we have np.arange(0, 10_000, dtype=np.int32)

2.

• [x] Do we need to more explicitly guide users about updating vector_add_gpu at the end of https://carpentries-incubator.github.io/lesson-gpu-programming/05-first_program/index.html I got myself into a bit of a pickle just copy-pasting code from the workshop notes into Jupyter and running them. I found I couldn't replicate the final Correct results! output at the end of Your First GPU Kernel as we don't actually explicitly show updating the kernel code in Python. Would it be worth adding an additional code cell updating vector_add_gpu with the new kernel before the steps setting more flexible grid and block size?

  # CUDA vector_add
  vector_add_gpu = cupy.RawKernel(r'''
  extern "C"
  __global__ void vector_add(const float * A, const float * B, float * C, const int size)
  {
   int item = (blockIdx.x * blockDim.x) + threadIdx.x;
   if ( item < size )
   {
      C[item] = A[item] + B[item];
   }
  }
  ''', "vector_add")
  

  3.

• [x] Couldn't get suggested answer for https://carpentries-incubator.github.io/lesson-gpu-programming/07-shared_memory_and_synchronization/index.html#challenge-use-shared-memory-to-speed-up-the-histogram to run When testing out the suggested answer to the above challenge I always got the following error

   CUDADriverError                           Traceback (most recent call last)
  Cell In[43], line 44
     42 # check correctness
     43 histogram(input_cpu, output_cpu)
  ---> 44 histogram_gpu(grid_size, block_size, (input_gpu, output_gpu))
     45 if np.allclose(output_cpu, output_gpu):
     46     print("Correct results!")
  
  File cupy/_core/raw.pyx:89, in cupy._core.raw.RawKernel.__call__()
  
  File cupy/_core/raw.pyx:96, in cupy._core.raw.RawKernel.kernel.__get__()
  
  File cupy/_core/raw.pyx:117, in cupy._core.raw.RawKernel._kernel()
  
  File cupy/cuda/function.pyx:276, in cupy.cuda.function.Module.get_function()
  
  File cupy/cuda/function.pyx:217, in cupy.cuda.function.Function.__init__()
  
  File cupy_backends/cuda/api/driver.pyx:226, in cupy_backends.cuda.api.driver.moduleGetFunction()
  
  File cupy_backends/cuda/api/driver.pyx:60, in cupy_backends.cuda.api.driver.check_status()
  
  CUDADriverError: CUDA_ERROR_NOT_FOUND: named symbol not found

  Is that to be expected given this isn't an ideal implementation?

[Sparrow0hawk]

Happy to suggest a PR adding in section for point 2 but would like to hear others thoughts!

[isazi]

* [ ]  Couldn't get suggested answer for https://carpentries-incubator.github.io/lesson-gpu-programming/07-shared_memory_and_synchronization/index.html#challenge-use-shared-memory-to-speed-up-the-histogram to run
  When testing out the suggested answer to the above challenge I always got the following error
  ```
    CUDADriverError                           Traceback (most recent call last)
  Cell In[43], line 44
      42 # check correctness
      43 histogram(input_cpu, output_cpu)
  ---> 44 histogram_gpu(grid_size, block_size, (input_gpu, output_gpu))
      45 if np.allclose(output_cpu, output_gpu):
      46     print("Correct results!")

  File cupy/_core/raw.pyx:89, in cupy._core.raw.RawKernel.__call__()

  File cupy/_core/raw.pyx:96, in cupy._core.raw.RawKernel.kernel.__get__()

  File cupy/_core/raw.pyx:117, in cupy._core.raw.RawKernel._kernel()

  File cupy/cuda/function.pyx:276, in cupy.cuda.function.Module.get_function()

  File cupy/cuda/function.pyx:217, in cupy.cuda.function.Function.__init__()

  File cupy_backends/cuda/api/driver.pyx:226, in cupy_backends.cuda.api.driver.moduleGetFunction()

  File cupy_backends/cuda/api/driver.pyx:60, in cupy_backends.cuda.api.driver.check_status()

  CUDADriverError: CUDA_ERROR_NOT_FOUND: named symbol not found
  ```
  Is that to be expected given this isn't an ideal implementation?

I believe while copy/pasting the code you also removed the extern "C" that was present. It was omitted in the solution because we put it before every kernel definition, but maybe should be there. Not sure.

[Sparrow0hawk]

It's exactly this reason, maybe it's worth making sure for every kernel definition in the text we include this just to be consistent?

[isazi]

* [ ]  Do we need to more explicitly guide users about updating `vector_add_gpu` at the end of https://carpentries-incubator.github.io/lesson-gpu-programming/05-first_program/index.html
  I got myself into a bit of a pickle just copy-pasting code from the workshop notes into Jupyter and running them. I found I couldn't replicate the final `Correct results!` output at the end of Your First GPU Kernel as we don't actually explicitly show updating the kernel code in Python. Would it be worth adding an additional code cell updating `vector_add_gpu` with the new kernel before the steps setting more flexible grid and block size?
  ```python
  # CUDA vector_add
  vector_add_gpu = cupy.RawKernel(r'''
  extern "C"
  __global__ void vector_add(const float * A, const float * B, float * C, const int size)
  {
     int item = (blockIdx.x * blockDim.x) + threadIdx.x;
     if ( item < size )
     {
        C[item] = A[item] + B[item];
     }
  }
  ''', "vector_add")
  ```

The kernel is modified in the "more work than necessary" challenge, but it is probably a good idea to show the complete code in a single notebook cell as done for the other chapters.

[isazi]

It's exactly this reason, maybe it's worth making sure for every kernel definition in the text we include this just to be consistent?

You are right, it is like that in other solutions so it was a mistake on my side to not add it to that one.

[Sparrow0hawk]

I'll take both these tasks to action and submit some PR's back.