time_random123.h (5774B)
1 /* 2 Copyright 2010-2011, D. E. Shaw Research. 3 All rights reserved. 4 5 Redistribution and use in source and binary forms, with or without 6 modification, are permitted provided that the following conditions are 7 met: 8 9 * Redistributions of source code must retain the above copyright 10 notice, this list of conditions, and the following disclaimer. 11 12 * Redistributions in binary form must reproduce the above copyright 13 notice, this list of conditions, and the following disclaimer in the 14 documentation and/or other materials provided with the distribution. 15 16 * Neither the name of D. E. Shaw Research nor the names of its 17 contributors may be used to endorse or promote products derived from 18 this software without specific prior written permission. 19 20 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS 21 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT 22 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR 23 A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT 24 OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, 25 SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT 26 LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, 27 DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY 28 THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT 29 (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE 30 OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. 31 */ 32 #ifndef TIME_RANDOM123_H__ 33 #define TIME_RANDOM123_H__ 1 34 35 /* 36 * This file contains the performance timing kernels for Random123 37 * RNGs and a few conventional PRNGs, which have been hacked into 38 * this framework. This code should NOT be considered as an 39 * example of using PRNGs. A few macros try to keep it 40 * cross-platform across C (serial or threads), CUDA and OpenCL. 41 * The TEST_TPL + include util_expandtpl trick is used to 42 * template-generate kernels for all the different RNGs across 43 * NxW. 44 */ 45 46 #include <Random123/philox.h> 47 #include <Random123/threefry.h> 48 49 #if defined(__OPENCL_VERSION__) 50 #define KERNEL __kernel 51 #define MEMTYPE __global 52 #define xprintf(x) 53 #elif defined(__CUDA_ARCH__) 54 #define xprintf(x) 55 #else 56 #define xprintf(x) printf x 57 #endif 58 59 #ifndef KERNEL 60 #define KERNEL 61 #endif 62 63 #ifndef MEMTYPE 64 #define MEMTYPE 65 #endif 66 67 /* LOOK_AT forces the compiler to actually produce the code that 68 computes the elements of A. Without it, optimizing compilers can 69 elide some or all of the computation of A (i.e., the RNG we're 70 trying to get timings for). There are many ways to do this. A 71 perhaps more natural way would be to keep a running sum of all the 72 results so far, but we found that gcc 4.5 and 4.6 could unroll that 73 code into a fully SSE-ized loop, which appears to be 74 unrepresentative of the kinds of optimizations that are possible in 75 practice. I.e., we could not find any "real" use of the output of 76 the RNG that permitted SSE-ization of the RNG. */ 77 #define LOOK_AT(A, I, N) do{ \ 78 if (N==4) if(R123_BUILTIN_EXPECT(!(A.v[N>2?3:0]^A.v[N>2?2:0]^A.v[N>1?1:0]^A.v[0]), 0)) ++I; \ 79 if (N==2) if(R123_BUILTIN_EXPECT(!(A.v[N>1?1:0]^A.v[0]), 0)) ++I; \ 80 if (N==1) if(R123_BUILTIN_EXPECT(!(A.v[0]), 0)) ++I; \ 81 }while(0) 82 83 84 /* Macro that will expand later into all the Random123 PRNGs for NxW_R */ 85 /* XXX AMDAPPSDK 2.4 seemed unhappy with the first arg being uint, but 86 was ok when it was changed to uint64_t. It's now back to unsigned 87 because that seems more correct and generic. Nobody's using 2.4 88 any more, are they? Note that this macro is expanded into CPU, 89 CUDA and OpenCL "kernels", so it has to be generic. */ 90 #define TEST_TPL(NAME, N, W, R) \ 91 KERNEL void test_##NAME##N##x##W##_##R(unsigned n, NAME##N##x##W##_ctr_t ctrinit, NAME##N##x##W##_ukey_t uk, MEMTYPE NAME##N##x##W##_ctr_t *ctr) \ 92 { \ 93 unsigned tid = get_global_id(0); \ 94 unsigned i; \ 95 NAME##N##x##W##_ctr_t c, v={{0}}; \ 96 NAME##N##x##W##_key_t k=NAME##N##x##W##keyinit(uk); \ 97 c = ctrinit; \ 98 if( R == NAME##N##x##W##_rounds ){ \ 99 for (i = 0; i < n; ++i) { \ 100 v = NAME##N##x##W(c, k); \ 101 LOOK_AT(v, i, N); \ 102 c.v[0]++; \ 103 } \ 104 }else { \ 105 for (i = 0; i < n; ++i) { \ 106 v = NAME##N##x##W##_R(R, c, k); \ 107 /*xprintf(("1: %s k[0] %lx c[0] %lx v[0] %lx\n", #NAME #N "x" #W "_" #R, (unsigned long) k.v[0], (unsigned long) c.v[0], (unsigned long) v.v[0]));*/ \ 108 /*if (c.v[0] == 0) printline_##NAME##N##x##W##_##R(k, c, &v, 1);*/ \ 109 LOOK_AT(v, i, N); \ 110 c.v[0]++; \ 111 } \ 112 } \ 113 ctr[tid] = v; \ 114 } 115 116 /* Now expand TEST_TPL for all the relevant RNGs */ 117 #include "util_expandtpl.h" 118 119 #endif /* TIME_RANDOM123_H__ */