Use statically-allocated __shared__ memory for __CUDA_ARCH__ < 200