API tools faq
paste
Advertisement
Guest User

opencl code

a guest
Sep 15th, 2015
143
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C 9.62 KB | None | 0 0
raw download clone embed print report
  1. #define EXP_TABLE_SIZE 1000
  2. #define MAX_EXP 6
  3. #define MAX_SENTENCE_LENGTH 1024
  4. #define MAX_CODE_LENGTH 40
  5. #define MAX_SENTENCE_NUM 6
  6. #define ALIGNMENT_FACTOR 32
  7. #define THREADS_PER_WORD 128
  8. #define BLOCK_SIZE 128
  9.  
  10. kernel void device_memset(global float * array, int size){
  11.     int idx = get_global_id(0);
  12.     if (idx < size)
  13.         array[idx] = 0;
  14. }
  15.  
  16.  
  17. void reduceInWarp(volatile local float * f, int idInWarp){
  18.  
  19.     for (unsigned int i=THREADS_PER_WORD /2; i>32; i>>=1) {
  20.         if (idInWarp < i) {
  21.             f[idInWarp] += f[idInWarp + i];
  22.         }
  23.         barrier(CLK_LOCAL_MEM_FENCE);
  24.     }
  25.     if (idInWarp < 32){
  26.         f[idInWarp] += f[idInWarp + 32];
  27.         f[idInWarp] += f[idInWarp + 16];
  28.         f[idInWarp] += f[idInWarp + 8];
  29.         f[idInWarp] += f[idInWarp + 4];
  30.         f[idInWarp] += f[idInWarp + 2];
  31.         f[idInWarp] += f[idInWarp + 1];
  32.     }
  33. }
  34.  
  35. void reduceInWarp64(volatile local float * f, int idInWarp){
  36.  
  37.     for (unsigned int i=THREADS_PER_WORD /2; i>64; i>>=1) {
  38.         if (idInWarp < i) {
  39.             f[idInWarp] += f[idInWarp + i];
  40.         }
  41.         barrier(CLK_LOCAL_MEM_FENCE);
  42.     }
  43.     if (idInWarp < 64){
  44.         f[idInWarp] += f[idInWarp + 64];
  45.         f[idInWarp] += f[idInWarp + 32];
  46.         f[idInWarp] += f[idInWarp + 16];
  47.         f[idInWarp] += f[idInWarp + 8];
  48.         f[idInWarp] += f[idInWarp + 4];
  49.         f[idInWarp] += f[idInWarp + 2];
  50.         f[idInWarp] += f[idInWarp + 1];
  51.     }
  52. }
  53.  
  54. kernel void device_cbow(int sentence_num, int layer1_size, int layer1_size_aligned,
  55.         int window, int negative, int table_size, int vocab_size,
  56.          global int * d_sen,  global int * d_table,
  57.          global float * d_syn0, global float *d_syn1neg,
  58.          global unsigned int * d_random,  global float * expTable, volatile local float * shared ){
  59.  
  60.  
  61.     int sentence_position = (get_local_id(0) / THREADS_PER_WORD) + (get_local_size(0) / THREADS_PER_WORD) * get_group_id(0);
  62.     int idInWarp = get_local_id(0) % THREADS_PER_WORD;
  63.  
  64.  
  65.     volatile local float * f = &shared [ (get_local_id(0) / THREADS_PER_WORD) * THREADS_PER_WORD];
  66.     volatile local float * neu1 = &shared [ BLOCK_SIZE + (get_local_id(0) / THREADS_PER_WORD) * layer1_size_aligned];
  67.     volatile local float * neu1e= & shared[BLOCK_SIZE + (get_local_size(0) / THREADS_PER_WORD) * layer1_size_aligned + (get_local_id(0) / THREADS_PER_WORD) * layer1_size_aligned];
  68.  
  69.     if (sentence_position < MAX_SENTENCE_LENGTH) {
  70.         unsigned int next_random = d_random[sentence_position];
  71.  
  72.         for (int sentence_idx = 0; sentence_idx < sentence_num; sentence_idx++){
  73.  
  74.             for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD) neu1[c] = 0;
  75.             for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD) neu1e[c] = 0;
  76.  
  77.  
  78.  
  79.             next_random = next_random * (unsigned int) 1664525 + 1013904223;
  80.             int b = next_random % window;
  81.             int word = d_sen[sentence_idx * MAX_SENTENCE_LENGTH + sentence_position];
  82.             // in -> hidden
  83.             int cw = 0;
  84.             for (int a = b; a < window * 2 + 1 - b; a++)
  85.                 if (a != window) {
  86.                     int w = sentence_position - window + a;
  87.                     if (w < 0)
  88.                         continue;
  89.                     if (w>= MAX_SENTENCE_LENGTH)
  90.                         continue;
  91.                     int last_word = d_sen[sentence_idx * MAX_SENTENCE_LENGTH + w];
  92.                     for (int c = idInWarp; c < layer1_size; c+= THREADS_PER_WORD)
  93.                         neu1[c] += d_syn0[c + last_word * layer1_size_aligned];
  94.  
  95.                     cw++;
  96.                 }
  97.            
  98.             if (cw) {
  99.                 for (int c = idInWarp; c < layer1_size; c+= THREADS_PER_WORD)
  100.                     neu1[c] /= cw;
  101.            
  102.             // NEGATIVE SAMPLING
  103.             int target, label;
  104.             float alpha =((global float *) &d_sen[MAX_SENTENCE_NUM * MAX_SENTENCE_LENGTH + sentence_idx])[0];
  105.  
  106.             if (negative > 0)
  107.  
  108.                 for (int d = 0; d < negative + 1; d++) {
  109.  
  110.  
  111.                     if (d == 0) {
  112.                         target = word;
  113.                         label = 1;
  114.                     } else {
  115.                         next_random = next_random * (unsigned int) 1664525
  116.                                 + 1013904223;
  117.                         target = d_table[(next_random) % table_size];
  118.                         if (target == 0)
  119.                             target = next_random % (vocab_size - 1) + 1;
  120.                         if (target == word)
  121.                             continue;
  122.                         label = 0;
  123.                     }
  124.                     int l2 = target * layer1_size_aligned;
  125.                     f[idInWarp] = 0;
  126.                
  127.                    
  128.                     for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD){
  129.                         f[idInWarp] += neu1[c] * d_syn1neg[c + l2];  
  130.                     }
  131.                     barrier(CLK_LOCAL_MEM_FENCE);
  132.                     // Do reduction here;
  133.                     for (unsigned int i=THREADS_PER_WORD /2; i>32; i>>=1) {
  134.                     if (idInWarp < i) {
  135.                         f[idInWarp] += f[idInWarp + i];
  136.                     }
  137.                     barrier(CLK_LOCAL_MEM_FENCE);
  138.                     }
  139.                     if (idInWarp < 32){
  140.                         f[idInWarp] += f[idInWarp + 32];
  141.                         f[idInWarp] += f[idInWarp + 16];
  142.                         f[idInWarp] += f[idInWarp + 8];
  143.                         f[idInWarp] += f[idInWarp + 4];
  144.                         f[idInWarp] += f[idInWarp + 2];
  145.                         f[idInWarp] += f[idInWarp + 1];
  146.                     }
  147.  
  148.                     barrier(CLK_LOCAL_MEM_FENCE);
  149.                    
  150.                     float g;
  151.                     if (f[0] > MAX_EXP)
  152.                         g = (label - 1) * alpha;
  153.                     else if (f[0] < -MAX_EXP)
  154.                         g = (label - 0) * alpha;
  155.                     else
  156.                         g = (label - expTable[(int) ((f[0] + MAX_EXP)
  157.                                     * (EXP_TABLE_SIZE / MAX_EXP / 2))]) * alpha;
  158.  
  159.                     //barrier(CLK_LOCAL_MEM_FENCE);
  160.                     for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD)
  161.                         neu1e[c] += g * d_syn1neg[c + l2];
  162.                     for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD)
  163.                         d_syn1neg[c + l2] += g * neu1[c];
  164.                    
  165.                 }
  166.             // hidden -> in
  167.             for (int a = b; a < window * 2 + 1 - b; a++)
  168.                 if (a != window) {
  169.                     int w = sentence_position - window + a;
  170.                     if (w < 0)
  171.                         continue;
  172.                     if (w >= MAX_SENTENCE_LENGTH)
  173.                         continue;
  174.                     int last_word = d_sen[sentence_idx * MAX_SENTENCE_LENGTH + w];
  175.  
  176.                     for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD)
  177.                         d_syn0[c + last_word * layer1_size_aligned] += neu1e[c];
  178.  
  179.                 }
  180.             }
  181.  
  182.         }// End for sentence_idx
  183.         // Update d_random
  184.         if (idInWarp == 0 ) d_random[sentence_position] = next_random;
  185.    
  186.     }
  187. }
  188.  
  189.  
  190. kernel void device_cbow64(int sentence_num, int layer1_size, int layer1_size_aligned,
  191.         int window, int negative, int table_size, int vocab_size,
  192.         global int * d_sen, global int * d_table,
  193.         global float * d_syn0, global float *d_syn1neg,
  194.         global unsigned int * d_random, __constant float * expTable, volatile local float * shared ){
  195.  
  196.  
  197.     int sentence_position = (get_local_id(0) / THREADS_PER_WORD) + (get_local_size(0) / THREADS_PER_WORD) * get_group_id(0);
  198.     int idInWarp = get_local_id(0) % THREADS_PER_WORD;
  199.  
  200.  
  201.     volatile local float * f = shared + (get_local_id(0) / THREADS_PER_WORD) * THREADS_PER_WORD;
  202.     volatile local float * neu1 = shared + BLOCK_SIZE + (get_local_id(0) / THREADS_PER_WORD) * layer1_size_aligned;
  203.     volatile local float * neu1e= shared + BLOCK_SIZE + (get_local_size(0) / THREADS_PER_WORD) * layer1_size_aligned + (get_local_id(0) / THREADS_PER_WORD) * layer1_size_aligned;
  204.  
  205.     if (sentence_position < MAX_SENTENCE_LENGTH) {
  206.         unsigned int next_random = d_random[sentence_position];
  207.  
  208.         for (int sentence_idx = 0; sentence_idx < sentence_num; sentence_idx++){
  209.  
  210.             for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD) neu1[c] = 0;
  211.             for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD) neu1e[c] = 0;
  212.  
  213.  
  214.  
  215.             next_random = next_random * (unsigned int) 1664525 + 1013904223;
  216.             int b = next_random % window;
  217.             int word = d_sen[sentence_idx * MAX_SENTENCE_LENGTH + sentence_position];
  218.             // in -> hidden
  219.             int cw = 0;
  220.             for (int a = b; a < window * 2 + 1 - b; a++)
  221.                 if (a != window) {
  222.                     int w = sentence_position - window + a;
  223.                     if (w < 0)
  224.                         continue;
  225.                     if (w>= MAX_SENTENCE_LENGTH)
  226.                         continue;
  227.                     int last_word = d_sen[sentence_idx * MAX_SENTENCE_LENGTH + w];
  228.                     for (int c = idInWarp; c < layer1_size; c+= THREADS_PER_WORD)
  229.                         neu1[c] += d_syn0[c + last_word * layer1_size_aligned];
  230.  
  231.                     cw++;
  232.                 }
  233.            
  234.             if (cw) {
  235.                 for (int c = idInWarp; c < layer1_size; c+= THREADS_PER_WORD)
  236.                     neu1[c] /= cw;
  237.            
  238.             // NEGATIVE SAMPLING
  239.             int target, label;
  240.             float alpha =*((global float *) &d_sen[MAX_SENTENCE_NUM * MAX_SENTENCE_LENGTH + sentence_idx]);
  241.  
  242.             if (negative > 0)
  243.  
  244.                 for (int d = 0; d < negative + 1; d++) {
  245.  
  246.  
  247.                     if (d == 0) {
  248.                         target = word;
  249.                         label = 1;
  250.                     } else {
  251.                         next_random = next_random * (unsigned int) 1664525
  252.                                 + 1013904223;
  253.                         target = d_table[(next_random) % table_size];
  254.                         if (target == 0)
  255.                             target = next_random % (vocab_size - 1) + 1;
  256.                         if (target == word)
  257.                             continue;
  258.                         label = 0;
  259.                     }
  260.                     int l2 = target * layer1_size_aligned;
  261.                     f[idInWarp] = 0;
  262.                
  263.                    
  264.                     for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD){
  265.                         f[idInWarp] += neu1[c] * d_syn1neg[c + l2];  
  266.                     }
  267.                     barrier(CLK_LOCAL_MEM_FENCE);
  268.                     // Do reduction here;
  269.                     reduceInWarp64(f, idInWarp);
  270.  
  271.                     barrier(CLK_LOCAL_MEM_FENCE);
  272.                    
  273.                     float g;
  274.                     if (f[0] > MAX_EXP)
  275.                         g = (label - 1) * alpha;
  276.                     else if (f[0] < -MAX_EXP)
  277.                         g = (label - 0) * alpha;
  278.                     else
  279.                         g = (label - expTable[(int) ((f[0] + MAX_EXP)
  280.                                     * (EXP_TABLE_SIZE / MAX_EXP / 2))]) * alpha;
  281.  
  282.                     //barrier(CLK_LOCAL_MEM_FENCE);
  283.                     for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD)
  284.                         neu1e[c] += g * d_syn1neg[c + l2];
  285.                     for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD)
  286.                         d_syn1neg[c + l2] += g * neu1[c];
  287.                    
  288.                 }
  289.             // hidden -> in
  290.             for (int a = b; a < window * 2 + 1 - b; a++)
  291.                 if (a != window) {
  292.                     int w = sentence_position - window + a;
  293.                     if (w < 0)
  294.                         continue;
  295.                     if (w >= MAX_SENTENCE_LENGTH)
  296.                         continue;
  297.                     int last_word = d_sen[sentence_idx * MAX_SENTENCE_LENGTH + w];
  298.  
  299.                     for (int c = idInWarp; c < layer1_size; c+=THREADS_PER_WORD)
  300.                         d_syn0[c + last_word * layer1_size_aligned] += neu1e[c];
  301.  
  302.                 }
  303.             }
  304.         }// End for sentence_idx
  305.         // Update d_random
  306.         if (idInWarp == 0 ) d_random[sentence_position] = next_random;
  307.     }
  308. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement
Public Pastes
Advertisement
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy.  OK, I Understand
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!