Implement a CUDA kernel that performs a 1D cross-correlation of a single-channel input signal of length N with a single-channel kernel of length kN. This is the operation deep learning frameworks call "convolution" (no kernel flip). Stride is 1 and there is no padding (valid).
For each output position i:
output[i]=a=0∑kN−1input[i+a]⋅kernel[a]The output has length N−kN+1. All three buffers are flat float32 device arrays.
Input: N = 5, kN = 3, input = [1, 2, 3, 4, 5], kernel = [1, 0, -1]
Output: [-2.0, -2.0, -2.0]
Each length-3 window contributes input[i]−input[i+2], which is −2 everywhere on this ramp.
Input: N = 4, kN = 2, input = [1, 2, 3, 4], kernel = [0.5, 0.5]
Output: [1.5, 2.5, 3.5]
A simple two-tap box average across each adjacent pair.
Each output position owns an independent length-kN window of input reads, so threads do not need to synchronize or share state.
Map blockIdx.x, threadIdx.x to the output index i and remember the output length is outN=N−kN+1, not N.
For small kN the inner loop is tight enough that the compiler unrolls it; accumulating into a local float avoids repeated global memory writes.
Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.
Implement a CUDA kernel that performs a 1D cross-correlation of a single-channel input signal of length N with a single-channel kernel of length kN. This is the operation deep learning frameworks call "convolution" (no kernel flip). Stride is 1 and there is no padding (valid).
For each output position i:
output[i]=a=0∑kN−1input[i+a]⋅kernel[a]The output has length N−kN+1. All three buffers are flat float32 device arrays.
Input: N = 5, kN = 3, input = [1, 2, 3, 4, 5], kernel = [1, 0, -1]
Output: [-2.0, -2.0, -2.0]
Each length-3 window contributes input[i]−input[i+2], which is −2 everywhere on this ramp.
Input: N = 4, kN = 2, input = [1, 2, 3, 4], kernel = [0.5, 0.5]
Output: [1.5, 2.5, 3.5]
A simple two-tap box average across each adjacent pair.
Each output position owns an independent length-kN window of input reads, so threads do not need to synchronize or share state.
Map blockIdx.x, threadIdx.x to the output index i and remember the output length is outN=N−kN+1, not N.
For small kN the inner loop is tight enough that the compiler unrolls it; accumulating into a local float avoids repeated global memory writes.
Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.