1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
|
#include <stdio.h>
#define SIZE 1024
#define THREADS_PER_BLOCK 32
#define PART_THREADS 1
#define NUM_BLOCKS 1
#define I_PREC 4
#define O_PREC 4
__global__ void vector_add(int* A, int* B, int* res)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
res[tid] = A[tid] + B[tid];
}
__global__ void digit_serial_mad(unsigned* i_buffer, unsigned* i_synapse, unsigned* result, unsigned* accum)
{
unsigned tid = threadIdx.x + blockIdx.x * blockDim.x;
unsigned buffer;
unsigned synapse;
if (tid < PART_THREADS)
{
buffer = i_buffer[tid];
synapse = i_synapse[tid];
}
asm("/*");
asm("CPTX_BEGIN");
asm("bsmad.s32 %0, %1, %2, %3, %4, %5, %6, %7, %8;" : "=r"(result[tid]) :
"r"(I_PREC), "r"(O_PREC), "r"(buffer), "r"(0), "r"(0), "r"(0), "r"(synapse), "r"(accum[tid]));
asm("CPTX_END");
asm("*/");
}
int main()
{
// host values
unsigned *buffer = (unsigned*)malloc(sizeof(unsigned));
unsigned *synapse = (unsigned*)malloc(sizeof(unsigned));
unsigned *result = (unsigned*)calloc(THREADS_PER_BLOCK, sizeof(unsigned));
unsigned *accum = (unsigned*)calloc(THREADS_PER_BLOCK, sizeof(unsigned));
// assign host values
*buffer = 0x5000003F;
*synapse = 0x00000002;
*accum = 0;
// device pointers
unsigned *d_buffer;
unsigned *d_synapse;
unsigned *d_result;
unsigned *d_accum;
// allocate device memory
cudaMalloc(&d_buffer, sizeof(unsigned));
cudaMalloc(&d_synapse, sizeof(unsigned));
cudaMalloc(&d_result, sizeof(unsigned));
cudaMalloc(&d_accum, sizeof(unsigned));
// copy data to device
cudaMemcpy(d_buffer, buffer, sizeof(unsigned), cudaMemcpyHostToDevice);
cudaMemcpy(d_synapse, synapse, sizeof(unsigned), cudaMemcpyHostToDevice);
cudaMemcpy(d_result, result, sizeof(unsigned) * THREADS_PER_BLOCK, cudaMemcpyHostToDevice);
cudaMemcpy(d_accum, accum, sizeof(unsigned) * THREADS_PER_BLOCK, cudaMemcpyHostToDevice);
// call kernel
digit_serial_mad<<<NUM_BLOCKS, THREADS_PER_BLOCK>>>(d_buffer, d_synapse, d_result, d_accum);
// copy data back to host
cudaMemcpy(result, d_result, sizeof(unsigned) * THREADS_PER_BLOCK, cudaMemcpyDeviceToHost);
// read out result
printf("Result: %#X\n", result[0]);
// clean up device memory
cudaFree(d_buffer);
cudaFree(d_synapse);
cudaFree(d_result);
cudaFree(d_accum);
// clean up host memory
free(buffer);
free(synapse);
free(result);
free(accum);
}
|
