Fix latency bug

1 files changed, 76 insertions, 0 deletions

diff --git a/bsmad_test/bsmad_test.cu b/bsmad_test/bsmad_test.cu
new file mode 100644
index 0000000..9be6e74
--- /dev/null
+++ b/bsmad_test/bsmad_test.cu
@@ -0,0 +1,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);
+}