5 ms·
When I write CUDA kernels I am very specific in optimizing for nVidia GPU features: * I have 1024 threads per block * I have 48KB of shared memory pe
by waynecochran 9y ago
When I write CUDA kernels I am very specific in optimizing for nVidia GPU features:
* I have 1024 threads per block
* I have 48KB of shared memory per block
* I have 32 threads per warp and need make sure that
global to local memory reads are coalesced.
* Think SIMD and avoid branching as much as possible
My kernels usually follow a specific pattern:
(1) Read global memory into local memory: making sure that
if thread i reads memory[n], then thread i+1
reads memory[n+1].
(2) __syncthreads().
(3) Do computation in in most thread balanced way possible.
This very specific pattern doesn't really work elsewhere. In fact, optimizing in this fashion and then porting to C++ and elsewhere loses the specific optimization. Programming in more general way loses all the things that makes the program fast. Anyway, I definitely going to look this over more.
- deleted 9y ago[deleted]
- AlphaSite 9y agoIt’s portable Cpp in the sense that it’s portable to AMD GPUs, rather than being exclusively comparable with Nvidias Cuda.
- dragontamer 9y agoAMD's "HCC" is C++ with templates added onto it. Unlike CUDA, HCC implements all relevant features in C++ Templates alone. If you simply do 64-thread per "warp" (AMD's groups are per-64) and 32KB LDS aka Shared memory per block, you would be able to write portable high-performance code between AMD GPUs and NVidia GPUs. AMD seems like its a bit behind with regards to GPGPU adoption. But AMD's hardware seems to be a good bit cheaper. You can get HBM2 models at ~$2000 from AMD (Firepro WX9100, which is Vega architecture) Although... as they say... hardware is cheap. I'm sure most datacenters will prefer the $8000 NVidia V100 instead, because there are more people using that hardware. In particular, its easier to get started with a V100 due to AWS and other cloud-compute offerings.