GPU Challenges / CUDA

Constant Memory Lookup Table: CUDA Exercise

hard Runs on a real NVIDIA T4 GPU in your browser

CUDA constant memory is a small (64 KB), cached, read-only space that is fastest when every thread in a warp reads the same address. This exercise uses it for a lookup table and shows how to fill it from the host with cudaMemcpyToSymbol.

The problem

Use CUDA constant memory to store a sine lookup table. Compute sin(x) for 1024 values using the lookup table instead of calling sin() directly. This demonstrates constant memory caching for read-only data accessed by all threads.

Use CUDA constant memory as a read-only lookup table so all threads can reuse cached sine values without recomputing sin() inside the kernel.

Your goal

  • Convert each input angle into a valid table index.
  • Read from __constant__ memory instead of recomputing the sine value.
  • Store the looked-up sine result into the output array.

How to approach it

  1. Declare the table at file scope as __constant__ float table[SIZE].
  2. Fill it once from the host with cudaMemcpyToSymbol before launching the kernel.
  3. In the kernel, convert each input to a table index and read the precomputed value instead of calling sin().
  4. Compare a few outputs with the CPU sin() to check the error is within tolerance.

Common mistakes

  • Using cudaMemcpy with a device pointer instead of cudaMemcpyToSymbol for constant memory.
  • Exceeding the 64 KB constant memory limit with a large table.
  • Expecting a speedup when threads in a warp read different addresses, which serialises constant-memory reads.

What success looks like

  • Printed sine values should closely match the CPU reference values.
  • Angles should map into the 0..359 lookup-table range safely.
  • The kernel should preserve the one-thread-per-output pattern.

Starter code

Start from this scaffold. Fill in the TODOs, then run it against the hidden tests.

constant-memory.cu
12345678910111213141516
#include <cuda_runtime.h>
#include <math.h>
#define TABLE_SIZE 360

__constant__ float sinTable[TABLE_SIZE];

// Kernel is launched as:  lookupSine<<<ceil(n/256), 256>>>(angles, results, n)
// gridDim.x  = ceil(n / 256)   blockDim.x = 256
// Thread ID:  int i = blockIdx.x * blockDim.x + threadIdx.x;
__global__ void lookupSine(const float* angles, float* results, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        // TODO: Convert angles[i] to integer index (0 to TABLE_SIZE-1)
        // TODO: Read sinTable[index] and write to results[i]
    }
}
Solve it in the editor →Free account, no GPU or CUDA install needed.

Concepts used

  • Constant memory usage
  • Lookup-table based computation
  • Read-only data shared across threads

Learn them first:

More CUDA challenges