結論だけ書くと、
・NVIDIA Performance Primitives (NPP) 9.1は、cannyがとてつもなく遅い。OpenCVが6msの条件で、NPPは86ms。もちろんデバイスメモリにコピーする時間は含まない。おそらくバグなので、そのうち速くなるかもしれない。
・OpenCVのUMat（OpenCL）は、モルフォロジー変換（erode / dilate）がとてつもなく遅い。OpenCLを使わないほうが倍くらい速い。
・もちろん両者をつなぎ合わせる方法はない。デバイス→ホスト→デバイスのコピーは気が遠くなるほど遅い。
・OpenCVのUMatのUMatを使いながら、モルフォロジー変換だけ自分でカーネルを書くには、C++のビルド環境が必要。
・必要なカーネルを全部自分で書くのでなければ、何をどうやってもC++のビルド環境が必要なので、Pythonで書く意味がそもそもない。
PythonでGPGPUが流行らないわけがよくわかった。NVIDIAは昔はNPPでPythonをサポートしていたらしいが、今はやめている。地獄めぐりがお好きな向きはどうぞ。
追記：まったく同じ引数（元画像の中身は異なる）で連続して2度呼び出すと、2度目は1msを切ってくる。スクラッチバッファをmallocした直後の1回だけ遅いらしい。キャッシュの関係か。
追記：CUDAを使うコードのプロファイリングにはcudaDeviceSynchronize()が必要。CUDAは（見た目はそうは見えないが）非同期で動く。
