Фамилии не знаю как искать, мат есть :-D

                // Note: I've tried a manual reduction and not using it because:
                // i ] I see no benefit. (slightly more expensive than the driver optimized atomic)
                // ii] huge chance that OCG will fuck up and reorder the LDS/STS due to
                //     not being able to infer data dependencies between LDS/STS reads/writes.
                //     (see bug 840564)

            //  printf("fuck STS, pq[0]=%f\n",__half2float(*((__half*)(&((*(int4*)(&smem_pool[0])).x)))));
            //  printf("fuck STS, pq[1]=%f\n",__half2float(*((__half*)(&((*(int4*)(&smem_pool[512])).x)))));
            //  printf("fuck STS, pq[2]=%f\n",__half2float(*((__half*)(&((*(int4*)(&smem_pool[1024])).x)))));
            //  printf("fuck STS, pq[3]=%f\n",__half2float(*((__half*)(&((*(int4*)(&smem_pool[1536])).x)))));

        # XXXgth: this is FUCKING BROKEN!!!

    // Send methods setting up the SM version, which DirectAModel on Kepler is
    // keenly interested in.  Hardware, like honey badger, don't give a shit.