looking at shmem sizes
looking at shmem allocations
doing __syncthreads() after each 2d slice of a filter doesn't seem to have a measurable speed penalty. cool.
naive shared memory cooperative loading for better coalescing in convolution. accelerates some layers, slows other layers down. not sure if it's feasible to get good conv performance with just one implementation that has to handle a variety of problem sizes.