1. #undef _GLIBCXX_ATOMIC_BUILTINS
  2. #undef _GLIBCXX_USE_INT128
  3. #include "cuda_code.h"
  5. void checkCUDAError(const char *msg)
  6. {
  7. cudaError_t err = cudaGetLastError();
  8. if (cudaSuccess != err) {
  9. printf("Cuda error: %s : %s.\n", msg, cudaGetErrorString(err));
  10. exit(EXIT_FAILURE);
  11. }
  12. }
  14. /// ACTIVATION
  15. __device__
  16. float Func(float number, FunctionType functionType)
  17. {
  18. switch (functionType) {
  20. case FT_BINARY_STEP:
  21. if (number > 0) {
  22. return 1;
  23. } else {
  24. return 0;
  25. }
  26. case FT_BIPOLAR_STEP:
  27. if (number > 0) {
  28. return 1;
  29. } else {
  30. return -1;
  31. }
  32. case SIGMOID:
  33. return 1.0f / (1.0f - exp(-number));
  34. case FT_BIPOLAR_SIGMOID:
  35. return -1.0f + (2.0f / (1.0f + exp(-number)));
  36. case FT_HYPERBOLIC_TANGENT:
  37. return tanh(number);
  38. case FT_IDENTITY:
  39. default:
  40. return number;
  41. }
  42. }
  44. __device__
  45. unsigned device_min(unsigned a, unsigned b)
  46. {
  47. if (a < b)
  48. return a;
  49. return b;
  50. }
  52. __global__
  53. void activation_float_kernel(float* results, float* thresholds, float* output, unsigned output_sz,
  54. FunctionType functionType)
  55. {
  56. int idx = blockIdx.x * blockDim.x + threadIdx.x;
  57. if (idx < output_sz) {
  58. output[idx] = Func(results[idx] - thresholds[idx], functionType);
  59. }
  60. }
  62. __global__
  63. void activation_bit_kernel(float* results, float* thresholds, unsigned* output, unsigned output_sz)
  64. {
  65. int idx = blockIdx.x * blockDim.x + threadIdx.x;
  66. unsigned offset = idx * BITS_PER_UNSIGNED;
  68. if (output_sz > offset) {
  70. unsigned toRead = device_min(BITS_PER_UNSIGNED, output_sz - offset);
  71. unsigned threadOutput = 0;
  72. unsigned mask = 0x80000000;
  74. for (unsigned i = 0; i < toRead; i++) {
  75. unsigned pos = offset + i;
  76. if (results[pos] - thresholds[pos] > 0) {
  77. threadOutput |= mask;
  78. } else {
  79. threadOutput &= ~mask;
  80. }
  81. mask >>= 1;
  82. }
  83. output[idx] = threadOutput;
  84. }
  85. }
  87. extern "C" void cuda_activation(void* output, unsigned size, BufferType bufferType, float* results,
  88. float* thresholds, FunctionType functionType, unsigned block_size)
  89. {
  90. unsigned grid_size;
  92. switch (bufferType) {
  93. case BT_BYTE:
  94. {
  95. std::string error = "cuda_activation is not implemented for BufferType BYTE.";
  96. throw error;
  97. }
  98. case BT_FLOAT:
  99. {
  100. grid_size = ((size - 1) / block_size) + 1;
  101. activation_float_kernel<<< grid_size, block_size >>>(results, thresholds, (float*)output, size, functionType);
  102. }
  103. break;
  104. case BT_BIT:
  105. case BT_SIGN:
  106. {
  107. grid_size = ((size - 1) / (block_size * BITS_PER_UNSIGNED)) + 1;
  108. activation_bit_kernel<<< grid_size, block_size >>>(results, thresholds, (unsigned*)output, size);
  109. }
  110. break;
  111. }
  112. checkCUDAError("activation");
  113. }
  115. // MEMORY MANAGEMENT
  117. extern "C" void* cuda_malloc(unsigned byteSize)
  118. {
  119. void* ptr;
  120. cudaMalloc((void**) &(ptr), byteSize);
  122. checkCUDAError("malloc");
  123. return ptr;
  124. }
  126. extern "C" void cuda_free(void* d_ptr)
  127. {
  128. cudaFree(d_ptr);
  129. checkCUDAError("free");
  130. }
  132. extern "C" void cuda_copyToDevice(void* d_dest, void* h_src, unsigned count)
  133. {
  134. cudaMemcpy(d_dest, h_src, count, cudaMemcpyHostToDevice);
  135. checkCUDAError("copyToDevice");
  136. }
  138. extern "C" void cuda_copyToHost(void* h_dest, void* d_src, unsigned count)
  139. {
  140. cudaMemcpy(h_dest, d_src, count, cudaMemcpyDeviceToHost);
  141. checkCUDAError("copyToHost");
  142. }
  144. // INITIALIZATION
  146. template <class bufferType>
  147. __global__
  148. void SetValueToAnArrayKernel(bufferType* data, unsigned size, bufferType value)
  149. {
  150. int idx = blockIdx.x * blockDim.x + threadIdx.x;
  151. if (idx < size)
  152. data[idx] = value;
  153. }
  155. extern "C" void cuda_setZero(void* data, unsigned byteSize, BufferType bufferType, unsigned block_size)
  156. {
  157. unsigned grid_size;
  158. unsigned size;
  160. switch (bufferType) {
  161. case BT_BYTE:
  162. size = byteSize / sizeof(unsigned char);
  163. grid_size = ((size - 1) / block_size) + 1;
  164. SetValueToAnArrayKernel<unsigned char><<< grid_size, block_size >>>((unsigned char*)data, size, (unsigned char)0);
  165. break;
  166. case BT_FLOAT:
  167. size = byteSize / sizeof(float);
  168. grid_size = ((size - 1) / block_size) + 1;
  169. SetValueToAnArrayKernel<float><<< grid_size, block_size >>>((float*)data, size, 0);
  170. break;
  171. case BT_BIT:
  172. case BT_SIGN:
  173. cudaMemset(data, 0, byteSize);
  174. break;
  175. }
  176. }
  178. // GENETIC OPERATORS
  180. template <class type>
  181. __global__
  182. void crossoverKernel(type* buffer1, type* buffer2, unsigned* bitBuffer, unsigned size)
  183. {
  184. unsigned weighPos = (blockIdx.x * blockDim.x * BITS_PER_UNSIGNED) + threadIdx.x;
  185. unsigned maxPosForThisBlock = device_min ( (blockIdx.x + 1) * blockDim.x * BITS_PER_UNSIGNED,
  186. size);
  187. unsigned bitsForTheThread, mask;
  188. if (weighPos < maxPosForThisBlock) {
  189. bitsForTheThread = bitBuffer[(blockIdx.x * blockDim.x) + threadIdx.x];
  190. mask = 0x80000000;
  191. }
  192. __syncthreads();
  193. while (weighPos < maxPosForThisBlock) {
  194. if (mask & bitsForTheThread) {
  195. type aux = buffer1[weighPos];
  196. buffer1[weighPos] = buffer2[weighPos];
  197. buffer2[weighPos] = aux;
  198. }
  199. weighPos += blockDim.x;
  200. mask >>= 1;
  201. }
  202. }
  204. extern "C" void cuda_crossover(void* buffer1, void* buffer2, unsigned* bitBuffer, unsigned size,
  205. BufferType bufferType, unsigned block_size)
  206. {
  207. unsigned grid_size = ((size - 1) / (block_size * BITS_PER_UNSIGNED)) + 1;
  209. switch (bufferType) {
  210. case BT_BYTE:
  211. crossoverKernel<unsigned char><<< grid_size, block_size >>>
  212. ((unsigned char*)buffer1, (unsigned char*)buffer2, (unsigned*)bitBuffer, size);
  214. break;
  215. case BT_FLOAT:
  216. crossoverKernel<float><<< grid_size, block_size >>>
  217. ((float*)buffer1, (float*)buffer2, (unsigned*)bitBuffer, size);
  218. break;
  219. case BT_BIT:
  220. case BT_SIGN:
  221. {
  222. std::string error = "cuda_crossover is not implemented for BufferType BIT nor SIGN.";
  223. throw error;
  224. }
  225. }
  226. }
  228. //TODO CU es necesario usar un kernel para esto ??
  229. __global__
  230. void resetFloatKernel(float* buffer, unsigned pos)
  231. {
  232. if (threadIdx.x == 0) {
  233. buffer[pos] = 0;
  234. }
  235. }
  237. __global__
  238. void resetByteKernel(unsigned char* buffer, unsigned pos)
  239. {
  240. if (threadIdx.x == 0) {
  241. buffer[pos] = 128;
  242. }
  243. }
  245. __global__
  246. void mutateFloatKernel(float* buffer, unsigned pos, float mutation)
  247. {
  248. if (threadIdx.x == 0) {
  249. buffer[pos] += mutation;
  250. }
  251. }
  253. __global__
  254. void mutateByteKernel(unsigned char* buffer, unsigned pos, int mutation)
  255. {
  256. if (threadIdx.x == 0) {
  257. int result = mutation + buffer[pos];
  258. if (result <= 0) {
  259. buffer[pos] = 0;
  260. } else if (result >= 255) {
  261. buffer[pos] = 255;
  262. } else {
  263. buffer[pos] = (unsigned char) result;
  264. }
  265. }
  266. }
  268. extern "C" void cuda_mutateWeigh(void* buffer, unsigned pos, float mutation, BufferType bufferType)
  269. {
  270. switch (bufferType) {
  271. case BT_BYTE:
  272. mutateByteKernel<<< 1, 8 >>>((unsigned char*)buffer, pos, (int)mutation);
  273. break;
  274. case BT_FLOAT:
  275. mutateFloatKernel<<< 1, 8 >>>((float*)buffer, pos, mutation);
  276. break;
  277. case BT_BIT:
  278. case BT_SIGN:
  279. {
  280. std::string error = "cuda_mutateWeigh is not implemented for BufferType BIT nor SIGN.";
  281. throw error;
  282. }
  283. }
  284. }
  286. extern "C" void cuda_resetWeigh(void* buffer, unsigned pos, BufferType bufferType)
  287. {
  288. switch (bufferType) {
  289. case BT_BYTE:
  290. resetByteKernel<<< 1, 8 >>>((unsigned char*)buffer, pos);
  291. break;
  292. case BT_FLOAT:
  293. resetFloatKernel<<< 1, 8 >>>((float*)buffer, pos);
  294. break;
  295. case BT_BIT:
  296. case BT_SIGN:
  297. {
  298. std::string error = "cuda_resetWeigh is not implemented for BufferType BIT nor SIGN.";
  299. throw error;
  300. }
  301. }
  302. }
  304. // CL_LAYER CALCULATION
  306. __global__
  307. void SumFloatsConnectionsKernel(float* inputs, unsigned input_size, unsigned output_size, float* weighs,
  308. float* results)
  309. {
  310. extern __shared__ float sdata[];
  312. unsigned outputNeuron = blockIdx.x * blockDim.x + threadIdx.x;
  313. unsigned weighsOffset = outputNeuron * input_size;
  314. float result = 0;
  316. unsigned pos = threadIdx.x;
  317. while (pos < input_size) {
  319. sdata[pos] = inputs[pos];
  320. pos += blockDim.x;
  321. }
  322. __syncthreads();
  324. if (outputNeuron < output_size) {
  326. //////////////////////////
  327. for (unsigned i = 0; i < input_size; i++) {
  328. result += sdata[i] * weighs[weighsOffset + i];
  329. //printf(" peso %f ", weighs[weighsOffset + i]);
  330. }
  331. /////TODO TR OTRA OPCION
  332. /* if (blockDim.x <= input_size){
  333. unsigned pos = tid;
  334. while (pos < input_size){
  335. result += sdata[pos] * weighs[weighsOffset + pos];
  336. ++pos;
  337. }
  338. pos = 0;
  339. while (pos < tid){
  340. result += sdata[pos] * weighs[weighsOffset + pos];
  341. ++pos;
  342. }
  343. } else {
  344. unsigned pos = tid;
  345. while (pos < input_size){
  346. result += sdata[pos] * weighs[weighsOffset + pos];
  347. ++pos;
  348. }
  349. unsigned newMax = device_min(tid, input_size);
  350. pos = 0;
  351. while (pos < newMax){
  352. result += sdata[pos] * weighs[weighsOffset + pos];
  353. ++pos;
  354. }
  355. }*/
  356. /////////////
  357. results[outputNeuron] += result;
  358. }
  359. }
  361. template <BufferType inputType>
  362. __global__
  363. void SumBitsConnectionsKernel(unsigned* inputs, unsigned input_size, unsigned output_size, unsigned char* weighs, float* results)
  364. {
  365. extern __shared__ unsigned shared_inputs[];
  367. unsigned tid = threadIdx.x;
  368. unsigned input_blocks_to_read = ((input_size - 1) / BITS_PER_UNSIGNED) + 1;
  369. unsigned readingLoops = ((input_blocks_to_read - 1) / blockDim.x) + 1;
  371. unsigned pos = tid;
  373. for (unsigned i=0; i < readingLoops; i++) {
  374. if (pos < input_blocks_to_read) {
  375. shared_inputs[pos] = inputs[pos];
  376. }
  377. pos += blockDim.x;
  378. }
  379. __syncthreads();
  381. unsigned outputNeuron = blockIdx.x*blockDim.x + threadIdx.x;
  382. if (outputNeuron < output_size) {
  384. float result = 0;
  385. unsigned weighsOffset = (outputNeuron * input_size);
  387. for (unsigned i=0; i < input_blocks_to_read; i++) {
  389. //TODO TCC check performance penalty (this is just for BT_SIGN)
  390. unsigned maxBits = device_min(BITS_PER_UNSIGNED, input_size - (i * BITS_PER_UNSIGNED));
  392. unsigned input_block = shared_inputs[i];
  393. unsigned mask = 0x80000000;
  394. for (unsigned j=0; j < maxBits; j++) {
  396. if (input_block & mask) {
  397. result += weighs[weighsOffset] - 128;
  398. } else {
  399. if (inputType == BT_SIGN) {
  400. result += 128 - weighs[weighsOffset];
  401. }
  402. }
  403. ++weighsOffset;
  404. mask >>= 1;
  405. }
  406. }
  407. results[outputNeuron] += result;
  408. }
  409. }
  411. __global__
  412. void SumFloatsInvertedConnectionsKernel(float* inputs, unsigned input_size, float* weighs, float* results,
  413. unsigned output_size)
  414. {
  415. extern __shared__ float sdata[];
  417. unsigned input_pos = threadIdx.x;
  418. while (input_pos < input_size) {
  420. sdata[input_pos] = inputs[input_pos];
  421. input_pos += blockDim.x;
  422. }
  423. __syncthreads();
  425. unsigned output_pos = blockIdx.x * blockDim.x + threadIdx.x;
  426. float result = 0;
  428. if (output_pos < output_size) {
  430. for (unsigned i = 0; i < input_size; i++) {
  431. result += sdata[i] * weighs[output_pos + (i * output_size)];
  432. }
  433. results[output_pos] += result;
  434. }
  435. }
  437. template <BufferType inputType>
  438. __global__
  439. void SumBitsInvertedConnectionsKernel(unsigned* inputs, unsigned input_size, unsigned output_size, unsigned char* weighs, float* results)
  440. {
  441. extern __shared__ unsigned shared_inputs[];
  443. unsigned tid = threadIdx.x;
  444. unsigned input_blocks_to_read = ((input_size - 1) / BITS_PER_UNSIGNED) + 1;
  445. unsigned readingLoops = ((input_blocks_to_read - 1) / blockDim.x) + 1;
  447. unsigned pos = tid;
  449. for (unsigned i=0; i < readingLoops; i++) {
  450. if (pos < input_blocks_to_read) {
  451. shared_inputs[pos] = inputs[pos];
  452. }
  453. pos += blockDim.x;
  454. }
  455. __syncthreads();
  457. unsigned outputNeuron = blockIdx.x*blockDim.x + threadIdx.x;
  458. if (outputNeuron < output_size) {
  460. float result = 0;
  462. for (unsigned i=0; i < input_blocks_to_read; i++) {
  464. //TODO TCC check performance penalty (this is just for BT_SIGN)
  465. unsigned maxBits = device_min(BITS_PER_UNSIGNED, input_size - (i * BITS_PER_UNSIGNED));
  467. unsigned weighsOffset = (i * BITS_PER_UNSIGNED * output_size) + outputNeuron;
  468. unsigned input_block = shared_inputs[i];
  469. unsigned mask = 0x80000000;
  470. for (unsigned j=0; j < maxBits; j++) {
  472. if (input_block & mask) {
  473. result += weighs[weighsOffset] - 128;
  474. } else {
  475. if (inputType == BT_SIGN) {
  476. result += 128 - weighs[weighsOffset];
  477. }
  478. }
  479. weighsOffset += output_size;
  480. mask >>= 1;
  481. }
  482. }
  483. results[outputNeuron] += result;
  484. }
  485. }
  487. extern "C" void cuda_inputCalculation(void* inputPtr, unsigned input_size, BufferType inputType,
  488. unsigned output_size, void* weighs, float* results, unsigned block_size)
  489. {
  490. unsigned grid_size = ((output_size - 1) / block_size) + 1;
  491. unsigned shared_mem_size;
  493. if (inputType == BT_BYTE) {
  494. std::string error = "cuda_inputCalculation is not implemented for BufferType BYTE as input.";
  495. throw error;
  496. } else if (inputType == BT_FLOAT) {
  497. if (input_size > 4032) {
  498. string error = "The maximum float input size is 4032.";
  499. throw error;
  500. }
  501. shared_mem_size = input_size * sizeof(float);
  503. SumFloatsConnectionsKernel<<< grid_size, block_size, shared_mem_size >>>((float*)inputPtr, input_size, output_size, (float*)weighs, results);
  504. } else {
  506. shared_mem_size =(((input_size - 1)/BITS_PER_UNSIGNED) + 1) * sizeof(unsigned);
  507. if (shared_mem_size > 16128) {
  508. //16128 * 8
  509. string error = "The maximum bit/sign input size is 129024.";
  510. throw error;
  511. }
  512. if (inputType == BT_BIT) {
  513. SumBitsConnectionsKernel<BT_BIT><<< grid_size, block_size, shared_mem_size >>>((unsigned*)inputPtr, input_size, output_size, (unsigned char*)weighs, results);
  514. } else {
  515. SumBitsConnectionsKernel<BT_SIGN><<< grid_size, block_size, shared_mem_size >>>((unsigned*)inputPtr, input_size, output_size, (unsigned char*)weighs, results);
  516. }
  517. }
  518. }
  520. extern "C" void cuda_inputCalculationInvertedMatrix(void* inputPtr, unsigned input_size, BufferType inputType,
  521. unsigned output_size, void* weighs, float* results,
  522. unsigned block_size)
  523. {
  524. unsigned grid_size = ((output_size - 1) / block_size) + 1;
  525. unsigned shared_mem_size;
  527. if (inputType == BT_BYTE) {
  528. std::string error = "cuda_inputCalculation is not implemented for BufferType BYTE as input.";
  529. throw error;
  530. } else if (inputType == BT_FLOAT) {
  531. while (input_size > CUDA_MAX_SHARED_FLOATS) {
  533. shared_mem_size = CUDA_MAX_SHARED_FLOATS * sizeof(float);
  534. SumFloatsInvertedConnectionsKernel<<< grid_size, block_size, shared_mem_size >>>((float*)inputPtr, CUDA_MAX_SHARED_FLOATS, (float*)weighs, results, output_size);
  535. inputPtr = (void*) ((float*) inputPtr + CUDA_MAX_SHARED_FLOATS);
  536. weighs = (void*) ((float*) weighs + (CUDA_MAX_SHARED_FLOATS * output_size));
  537. input_size -= CUDA_MAX_SHARED_FLOATS;
  538. }
  539. shared_mem_size = input_size * sizeof(float);
  540. SumFloatsInvertedConnectionsKernel <<< grid_size, block_size, shared_mem_size >>>((float*)inputPtr, input_size, (float*)weighs, results, output_size);
  541. } else {
  542. //TODO TCC esta parte no funciona bien
  543. while (input_size > CUDA_MAX_SHARED_BITS) {
  545. shared_mem_size = CUDA_MAX_SHARED_FLOATS * sizeof(unsigned);
  546. // TODO TCC probar sin emulación
  547. // printf("grid_size %d, block_size %d, shared_mem_size %d \n", grid_size, block_size, shared_mem_size);
  548. if (inputType == BT_BIT) {
  549. SumBitsInvertedConnectionsKernel<BT_BIT><<< grid_size, block_size, shared_mem_size >>>((unsigned*)inputPtr, CUDA_MAX_SHARED_BITS, output_size, (unsigned char*)weighs, results);
  550. } else {
  551. SumBitsInvertedConnectionsKernel<BT_SIGN><<< grid_size, block_size, shared_mem_size >>>((unsigned*)inputPtr, CUDA_MAX_SHARED_BITS, output_size, (unsigned char*)weighs, results);
  552. }
  553. inputPtr = (void*)((float*)inputPtr + CUDA_MAX_SHARED_FLOATS);
  554. weighs = (void*)((float*)weighs + (CUDA_MAX_SHARED_BITS * output_size));
  555. input_size -= CUDA_MAX_SHARED_BITS;
  556. }
  557. shared_mem_size =(((input_size - 1)/BITS_PER_UNSIGNED) + 1) * sizeof(unsigned);
  558. // TODO TCC probar sin emulación
  559. //printf("grid_size %d, block_size %d, shared_mem_size %d \n", grid_size, block_size, shared_mem_size);
  560. if (inputType == BT_BIT) {
  561. SumBitsInvertedConnectionsKernel<BT_BIT><<< grid_size, block_size, shared_mem_size >>>((unsigned*)inputPtr, input_size, output_size, (unsigned char*)weighs, results);
  562. } else {
  563. SumBitsInvertedConnectionsKernel<BT_SIGN><<< grid_size, block_size, shared_mem_size >>>((unsigned*)inputPtr, input_size, output_size, (unsigned char*)weighs, results);
  564. }
  565. }
  566. }
  568. template <unsigned int blockSize, BufferType inputType>
  569. __global__
  570. void SumConnectionsKernel(void* inputPtr, unsigned input_size, unsigned output_size, void* weighs, float* results)
  571. {
  572. extern __shared__ float sdata[];
  574. unsigned weighsOffset = (blockIdx.x * input_size);
  576. float result = 0;
  577. unsigned i = threadIdx.x;
  579. if (inputType == BT_FLOAT) {
  580. while (i < input_size) {
  581. result += ((float*)inputPtr)[i] * ((float*)weighs)[weighsOffset + i];
  582. i += blockDim.x;
  583. }
  584. } else {
  585. weighsOffset += threadIdx.x * BITS_PER_UNSIGNED;
  587. unsigned input_blocks_to_read = ((input_size - 1) / BITS_PER_UNSIGNED) + 1;
  588. while (i < input_blocks_to_read) {
  590. //TODO TCC check performance penalty (this is just for BT_SIGN)
  591. unsigned maxBits = device_min(BITS_PER_UNSIGNED, input_size - (i * BITS_PER_UNSIGNED));
  593. unsigned mask = 0x80000000;
  594. unsigned currentInput = ((unsigned*)inputPtr)[i];
  596. for (unsigned j=0; j < maxBits; j++) {
  598. if (currentInput & mask) {
  599. result += ((unsigned char*)weighs)[weighsOffset + j] - 128;
  600. } else {
  601. if (inputType == BT_SIGN) {
  602. result -= ((unsigned char*)weighs)[weighsOffset + j] - 128;
  603. }
  604. }
  605. mask >>= 1;
  606. }
  607. i += blockSize;
  608. weighsOffset += blockDim.x * BITS_PER_UNSIGNED;
  609. }
  610. }
  612. unsigned tid = threadIdx.x;
  613. sdata[tid] = result;
  614. __syncthreads();
  616. if (blockSize >= 512) {if (tid < 256) {sdata[tid] += sdata[tid + 256];}__syncthreads();}
  617. if (blockSize >= 256) {if (tid < 128) {sdata[tid] += sdata[tid + 128];}__syncthreads();}
  618. if (blockSize >= 128) {if (tid < 64) {sdata[tid] += sdata[tid + 64];}__syncthreads();}
  620. #if __DEVICE_EMULATION__
  621. if (blockSize >= 64) {if (tid < 32) {sdata[tid] += sdata[tid + 32];}__syncthreads();}
  622. if (blockSize >= 32) {if (tid < 16) {sdata[tid] += sdata[tid + 16];}__syncthreads();}
  623. if (blockSize >= 16) {if (tid < 8) {sdata[tid] += sdata[tid + 8];}__syncthreads();}
  624. if (blockSize >= 8) {if (tid < 4) {sdata[tid] += sdata[tid + 4];}__syncthreads();}
  625. if (blockSize >= 4) {if (tid < 2) {sdata[tid] += sdata[tid + 2];}__syncthreads();}
  626. if (blockSize >= 2) {if (tid < 1) {sdata[tid] += sdata[tid + 1];}__syncthreads();}
  627. #else
  628. if (tid < 32) {
  629. if (blockSize >= 64) sdata[tid] += sdata[tid + 32];
  630. if (blockSize >= 32) sdata[tid] += sdata[tid + 16];
  631. if (blockSize >= 16) sdata[tid] += sdata[tid + 8];
  632. if (blockSize >= 8) sdata[tid] += sdata[tid + 4];
  633. if (blockSize >= 4) sdata[tid] += sdata[tid + 2];
  634. if (blockSize >= 2) sdata[tid] += sdata[tid + 1];
  635. }
  636. #endif
  637. if (tid == 0) {
  638. results[blockIdx.x] += sdata[0];
  639. }
  640. }
  642. extern "C" void cuda_inputCalculationReduction(void* inputPtr, unsigned input_size, BufferType inputType,
  643. unsigned output_size, void* weighs, float* results,
  644. unsigned block_size)
  645. {
  646. unsigned grid_size = output_size;
  647. unsigned shared_mem_size = block_size * sizeof(float);
  649. if (inputType == BT_BYTE) {
  650. std::string error = "cuda_inputCalculation is not implemented for BufferType BYTE as input.";
  651. throw error;
  652. } else if (inputType == BT_FLOAT) {
  653. switch (block_size) {
  654. case 512:
  655. SumConnectionsKernel<512, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  656. case 256:
  657. SumConnectionsKernel<256, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  658. case 128:
  659. SumConnectionsKernel<128, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  660. case 64:
  661. SumConnectionsKernel< 64, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  662. case 32:
  663. SumConnectionsKernel< 32, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  664. case 16:
  665. SumConnectionsKernel< 16, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  666. case 8:
  667. SumConnectionsKernel< 8, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  668. case 4:
  669. SumConnectionsKernel< 4, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  670. case 2:
  671. SumConnectionsKernel< 2, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  672. case 1:
  673. SumConnectionsKernel< 1, BT_FLOAT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  674. }
  675. } else if (inputType == BT_BIT) {
  676. switch (block_size) {
  677. case 512:
  678. SumConnectionsKernel<512, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  679. case 256:
  680. SumConnectionsKernel<256, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  681. case 128:
  682. SumConnectionsKernel<128, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  683. case 64:
  684. SumConnectionsKernel< 64, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  685. case 32:
  686. SumConnectionsKernel< 32, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  687. case 16:
  688. SumConnectionsKernel< 16, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  689. case 8:
  690. SumConnectionsKernel< 8, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  691. case 4:
  692. SumConnectionsKernel< 4, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  693. case 2:
  694. SumConnectionsKernel< 2, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  695. case 1:
  696. SumConnectionsKernel< 1, BT_BIT><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  697. }
  698. } else {
  699. switch (block_size) {
  700. case 512:
  701. SumConnectionsKernel<512, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  702. case 256:
  703. SumConnectionsKernel<256, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  704. case 128:
  705. SumConnectionsKernel<128, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  706. case 64:
  707. SumConnectionsKernel< 64, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  708. case 32:
  709. SumConnectionsKernel< 32, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  710. case 16:
  711. SumConnectionsKernel< 16, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  712. case 8:
  713. SumConnectionsKernel< 8, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  714. case 4:
  715. SumConnectionsKernel< 4, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  716. case 2:
  717. SumConnectionsKernel< 2, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  718. case 1:
  719. SumConnectionsKernel< 1, BT_SIGN><<< grid_size, block_size, shared_mem_size >>>(inputPtr, input_size, output_size, weighs, results); break;
  720. }
  721. }
  722. checkCUDAError("cuda_inputCalculation2");
  723. }