結論だけ書くと、
・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は(見た目はそうは見えないが)非同期で動く。