diff options
Diffstat (limited to 'bsmad_test/bsmad_test.cu')
| -rw-r--r-- | bsmad_test/bsmad_test.cu | 76 |
1 files changed, 0 insertions, 76 deletions
diff --git a/bsmad_test/bsmad_test.cu b/bsmad_test/bsmad_test.cu deleted file mode 100644 index 9be6e74..0000000 --- a/bsmad_test/bsmad_test.cu +++ /dev/null @@ -1,76 +0,0 @@ -#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); -} |
