diff --git a/hpvm/projects/hpvm-tensor-rt/dnn_sources/src/unit_tests.cc b/hpvm/projects/hpvm-tensor-rt/dnn_sources/src/unit_tests.cc
index 3b08755172973d63132bcd1c5b19d9e58ec38611..6793cd79f19bfe9fd192b2e6199323bb28940aa3 100644
--- a/hpvm/projects/hpvm-tensor-rt/dnn_sources/src/unit_tests.cc
+++ b/hpvm/projects/hpvm-tensor-rt/dnn_sources/src/unit_tests.cc
@@ -231,24 +231,6 @@ void testTensorAdd(){
printTensorValues(x2);
}
-
-void testTensorError(){
-
- // Tensor add with equal dimensions
- void* x = create4DTensor(CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, 2, 2, 2, 128);
- fillTensorWithOnes(x);
-
- Tensor* x_tensor = (Tensor*) x;
- float* data_arr = (float*) x_tensor->host_data;
- for(int i = 0; i < x_tensor->num_elems; i++){
- data_arr[i] = 0.2;
- }
-
- tensorAddError(x, 3);
- printTensorValues(x);
-}
-
-
void testTensorConv(){
void* input = create4DTensor(CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, 1, 2, 4, 4);
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/error.h b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/error.h
index a3d51141acd9e45d3231689a39f43e97fbeb0a9f..8c9a711c8a8355eb7e0240cc6ed15b5c7ebd23c9 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/error.h
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/error.h
@@ -45,21 +45,6 @@ __global__ void vecConstDiv(float *A, float div_factor, int n);
__global__ void vecMul(float *A, float *B, int n);
-/**** ERROR injecion routines ******/
-void initRandValues(Tensor *bias, int error_scale);
-
-void initRandValues2(Tensor *bias, int error_scale);
-
-void *addBitError(void *x_ptr, int error_scale);
-
-void randomCeilAndFloor(float *x, size_t num_elems);
-
-// Routine for Adding RoundOff Errors
-void *addRoundError(void *x_ptr, int error_scale);
-
-// Routine for Adding Gaussian Error
-void *addGaussianError(void *x_ptr, int error_scale);
-
void initPromiseRandValues(Tensor *bias, int error_scale);
// NOTE: Assumption is that x_ptr is FP32 tensor - doesn't work with FP16
@@ -72,8 +57,6 @@ __global__ void quantizeAndClip(float *A, int n, float mul_factor, float min,
__global__ void quantizeElem(float *A, int n, float mul_factor, float min);
void *quantizeTensorPromise(void *input_ptr, float min, float max);
-
-void *tensorAddError(void *x_ptr, int error_scale);
}
#endif
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/global_data.h b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/global_data.h
index f859b83e94ecc2b4e103792b977279613f119d71..49c1725336ab4242ba4ed852a10ba3cde0d1c5d7 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/global_data.h
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/global_data.h
@@ -20,14 +20,8 @@
#include "approx_knob_utils.h"
#include "tensor.h"
-#define ERROR_INJECTION_ENABLED 0
#define PROMISE_MODE 1
-#ifdef NO_INJECTION
-#undef ERROR_INJECTION_ENABLED
-#endif
-
-//#define ERROR_INJECTION_ENABLED 1
/* Data declarations */
extern cudnnHandle_t cudnnHandle;
extern cublasHandle_t cublasHandle;
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_cpu_runtime.h b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_cpu_runtime.h
index 4bb703bbd2596980fb4d930b36aaa749c7144044..f8b722ca38bc2ae7065da365585ca495001038d7 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_cpu_runtime.h
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_cpu_runtime.h
@@ -81,42 +81,4 @@ void* tensorConvCutlassCPU(void* input_ptr, void* filter_ptr,
}
-/*
-void dummyFunction(){
-
- void* initRT = (void*) &llvm_hpvm_initTensorRt;
- void* cleanRT = (void*) &llvm_hpvm_cleanupTensorRt;
- void* request_tensorPtr = (void*) &hpvm_request_tensor;
- void* startProf = (void*) &startProfiling;
- void* stopProf = (void*) &stopProfiling;
- void* create2Dptr = (void*) &create2DTensor;
- void* create3Dptr = (void*) &create3DTensor;
- void* create4Dptr = (void*) &create4DTensor;
- void* initTensorPtr = (void*) &initTensorData;
- void* tensorSplitPtr = (void*) &tensorSplit;
- void* tensorConcatPtr = (void*) &tensorConcat;
- void* tensorConvPtr = (void*) &tensorConvolution;
- void* tensorHConvPtr = (void*) &tensorHalfConvolution;
- void* tensorPoolPtr = (void*) &tensorPooling;
- void* tensorHalfPoolPtr = (void*) &tensorHalfPooling;
- void* tensorLRNPtr = (void*) &tensorLRN;
- void* tensorGemmPr = (void*) &tensorGemm;
- void* tensorGemmCPUPtr = (void*) &tensorGemmCPU;
- void* tensorGemmGPUPtr = (void*) &tensorGemmGPU;
- void* tensorHgemmPtr = (void*) &tensorHalfGemm;
- void* tensorGemmBiasPtr = (void*) &tensorGemmBias;
- void* tensorAddPtr = (void*) &tensorAdd;
- void* tensorHalfAddPtr = (void*) &tensorHalfAdd;
- void* tensorReluPtr = (void*) &tensorRelu;
- //FIXME: --void* tensorHalfReluPtr = (void*) &tensorHalfRelu;
- void* tensorRelu2Ptr = (void*) &tensorRelu2;
- void* tensorHalfRelu2Ptr = (void*) &tensorHalfRelu2;
- void* tensorTanhPtr = (void*) &tensorTanh;
- void* tensorHalfTanhPtr = (void*) &tensorHalfTanh;
- void* tensorSoftmaxPtr = (void*) &tensorSoftmax;
- void* tensorAddErrorPtr = (void*) &tensorAddError;
-}
-*/
-
-
#endif
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_runtime.cc b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_runtime.cc
index 0de2808221adfb122860a031eea4ed8c89d6e2ba..083d733b14d4f335f4365503e5ecdb5c3de0a2fb 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_runtime.cc
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_runtime.cc
@@ -69,37 +69,6 @@ void *tensorAdd(void *x, void *bias);
void *tensorRelu(void *input);
// NOTE: In-place operation
void *tensorSoftmax(void *input);
-
-/* Error injection API - used for accuracy tuning */
-void *tensorAddError(void *x_ptr);
-}
-
-void emptyFunction() {
-
- void *initRT = (void *)&llvm_hpvm_initTensorRt;
- void *cleanRT = (void *)&llvm_hpvm_cleanupTensorRt;
- void *request_tensorPtr = (void *)&hpvm_request_tensor;
- void *startProf = (void *)&startProfiling;
- void *stopProf = (void *)&stopProfiling;
- void *create2Dptr = (void *)&create2DTensor;
- void *create3Dptr = (void *)&create3DTensor;
- void *create4Dptr = (void *)&create4DTensor;
- void *initTensorPtr = (void *)&initTensorData;
- void *tensorSplitPtr = (void *)&tensorSplit;
- void *tensorConcatPtr = (void *)&tensorConcat;
- void *tensorConvPtr = (void *)&tensorConvolution;
- void *tensorHConvPtr = (void *)&tensorHConvolution;
- void *tensorPoolPtr = (void *)&tensorPooling;
- void *tensorLRNPtr = (void *)&tensorLRN;
- void *tensorGemmPr = (void *)&tensorGemm;
- void *tensorGemmCPUPtr = (void *)&tensorGemmCPU;
- void *tensorGemmGPUPtr = (void *)&tensorGemmGPU;
- void *tensorHgemmPtr = (void *)&tensorHgemm;
- void *tensorGemmBiasPtr = (void *)&tensorGemmBias;
- void *tensorAddPtr = (void *)&tensorAdd;
- void *tensorReluPtr = (void *)&tensorRelu;
- void *tensorSoftmaxPtr = (void *)&tensorSoftmax;
- void *tensorAddErrorPtr = (void *)&tensorAddError;
}
#endif
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_runtime.h b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_runtime.h
index eaa155d192cccf13bc683569253e0696fbccd960..f05dab738bbfab51e21673bdf76b81596fc3b49b 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_runtime.h
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_runtime.h
@@ -109,13 +109,6 @@ void *tensorBatchNorm(void *input_ptr, void *gamma_ptr, void *beta_ptr,
void *tensorHalfBatchNorm(void *input_ptr, void *gamma_ptr, void *beta_ptr,
void *mean_ptr, void *variance_ptr, double epsilon);
-/* Error injection API - used for accuracy tuning */
-void *tensorAddError(void *x_ptr, int error_scale);
-
-void *tensorGemmModel(void *lhs, void *rhs);
-
-/*** Error Injection API End **/
-
/**** PROMISE API *****/
/*************
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_signatures.cc b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_signatures.cc
index 73fa9df5b0d5de81350b30064d54673696bf374e..b3abdc0ce48a507a250e7f82fc6ff7729dff3e9f 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_signatures.cc
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/tensor_signatures.cc
@@ -51,7 +51,6 @@ void dummyFunction() {
void *tensorHalfTanhPtr = (void *)&tensorHalfTanh;
void *tensorSoftmaxPtr = (void *)&tensorSoftmax;
void *tensorBatchNormPtr = (void *)&tensorBatchNorm;
- void *tensorAddErrorPtr = (void *)&tensorAddError;
void *ConvLayer = (void *)&ConvLayer_PROMISE;
void *FCLayer = (void *)&FCLayer_PROMISE;
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/error.cu b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/error.cu
index 7a700b435efe464153fbba7997662c7dfa970385..4afed4c287e5c282fd6a4f43f7c4231e6b558fb4 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/error.cu
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/error.cu
@@ -548,307 +548,6 @@ __global__ void vecMul(float* A, float* B, int n){
B[id] = A[id] * B[id];
}
-
-/**** ERROR injecion routines ******/
-
-void initRandValues(Tensor* bias, int error_scale){
-
- float scaling_values[20];
-
- // FIXIT: Error knob 0 should be 0 zero
- scaling_values[0] = 0.000;
- scaling_values[1] = 0.0005;
- scaling_values[2] = 0.03;
- scaling_values[3] = 0.06;
- scaling_values[4] = 0.08;
- scaling_values[5] = 0.105;
- scaling_values[6] = 0.134;
- scaling_values[7] = 0.16;
- scaling_values[8] = 0.2;
- scaling_values[9] = 0.23;
- scaling_values[10] = 0.26;
- scaling_values[11] = 0.3;
- scaling_values[12] = 0.35;
- scaling_values[13] = 0.4;
- scaling_values[14] = 0.45;
- scaling_values[15] = 0.55;
- scaling_values[16] = 0.65;
- scaling_values[17] = 0.7;
- scaling_values[18] = 0.8;
- scaling_values[19] = 0.9;
-
-
- curandGenerator_t gen;
-
- struct timespec ts;
-
- if(timespec_get(&ts, TIME_UTC) == 0){
- printf("crashed \n");
- abort();
- }
-
- curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT);
-
- curandSetPseudoRandomGeneratorSeed(gen, ts.tv_nsec^ts.tv_sec);
-
- curandGenerateNormal(gen, (float*) bias->gpu_data, bias->num_elems, 0.0, 1.0 * scaling_values[error_scale]);
-
-}
-
-
-
-void initRandValues2(Tensor* bias, int error_scale){
-
- float scaling_values[20];
-
- // FIXIT: Error knob 0 should be 0 zero
- scaling_values[0] = 0.000;
- scaling_values[1] = 0.0005;
- scaling_values[2] = 0.0008;
- scaling_values[3] = 0.001;
- scaling_values[4] = 0.005;
- scaling_values[5] = 0.01;
- scaling_values[6] = 0.02;
- scaling_values[7] = 0.03;
- scaling_values[8] = 0.04;
- scaling_values[9] = 0.05;
- scaling_values[10] = 0.06;
- scaling_values[11] = 0.08;
- scaling_values[12] = 0.1;
- scaling_values[13] = 0.12;
- scaling_values[14] = 0.15;
- scaling_values[15] = 0.2;
- scaling_values[16] = 0.55;
- scaling_values[17] = 0.6;
- scaling_values[18] = 0.65;
- scaling_values[19] = 0.7;
-
-
- curandGenerator_t gen;
-
- struct timespec ts;
-
- if(timespec_get(&ts, TIME_UTC) == 0){
- printf("crashed \n");
- abort();
- }
-
- curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT);
-
- curandSetPseudoRandomGeneratorSeed(gen, ts.tv_nsec^ts.tv_sec);
-
- curandGenerateNormal(gen, (float*) bias->gpu_data, bias->num_elems, 0.0, 1.0 * scaling_values[error_scale]);
-
-}
-
-
-void* addBitError(void* x_ptr, int error_scale){
-
- if(error_scale > 6 || error_scale < 0){
- ERROR("Error Scale out of bounds \n");
- }
-
- INFO("*** TensorBitError \n");
- profileEvent("tensorBitError");
-
- Tensor* x = (Tensor*) x_ptr;
-
- size_t* dim_sizes = x->dims.dim_sizes;
- Tensor* x_original = (Tensor*) create4DTensor(x->data_type, x->data_format,
- dim_sizes[0], dim_sizes[1],
- dim_sizes[2], dim_sizes[3]);
-
- // Copying x data into x_original - for computing Norms
- tensorCopy(x, x_original);
-
- // Quadratic Error
- float freq_factors[6];
- freq_factors[0] = 0.1;
- freq_factors[1] = 0.2;
- freq_factors[2] = 0.4;
- freq_factors[3] = 0.6;
- freq_factors[4] = 0.8;
- freq_factors[5] = 1.0;
-
- float error_freq = freq_factors[error_scale];
-
- deviceToHostCopy(x);
-
- unsigned char* data_arr = reinterpret_cast<unsigned char*>(x->host_data);
- // FIXIT: Need to be careful about floating point datatype assumptions
- long int total_bytes = x->size_in_bytes;
- long int error_iterations = total_bytes * 0.01 * error_freq;
- INFO("total_bytes = %lu, error_iterations = %lu \n", total_bytes, error_iterations);
-
- srand(time(NULL));
-
- for(int i = 0; i < error_iterations; i++){
- // FIXIT: The rand() is only specific to int - need long
- long int index = rand() % total_bytes;
- int N = 5; // The operation below flips the Nth bit
- unsigned char fil = 1UL << N;
- unsigned char val = data_arr[index];
- char flipped = val^fil;
- data_arr[i] = flipped;
- }
-
-
- Norm_t* norms = calculateNorms2(x, x_original);
-
-
- profileEvent("tensorBitError_end", true);
-
- return (void*) norms;
-
-}
-
-
-void randomCeilAndFloor(float* x, size_t num_elems){
-
- INFO("randomCeilAndFloor\n");
-
- std::random_device rd;
- std::mt19937 mt(rd());
- std::normal_distribution<float> distribution(0.0, 1.0);
-
- for(size_t i = 0; i < num_elems; i++){
- float rand_num = distribution(mt);
- int val = abs(((int) rand_num) % 2);
- if(val == 0)
- x[i] = floor(x[i]);
- else if(val == 1)
- x[i] = ceil(x[i]);
- }
-
-}
-
-// Routine for Adding RoundOff Errors
-void* addRoundError(void* x_ptr, int error_scale){
-
- if(error_scale > 11 || error_scale < 0){
- ERROR("Error Scale out of bounds \n");
- }
-
- INFO("*** TensorRoundError \n");
- profileEvent("tensorRoundError");
-
- Tensor* x = (Tensor*) x_ptr;
-
- size_t* dim_sizes = x->dims.dim_sizes;
- Tensor* x_original = (Tensor*) create4DTensor(x->data_type, x->data_format,
- dim_sizes[0], dim_sizes[1],
- dim_sizes[2], dim_sizes[3]);
-
- // Copying x data into x_original - for computing Norms
- tensorCopy(x, x_original);
-
- float round_factors[12];
- round_factors[0] = 1000000; // FIXIT: This should be zero error
- round_factors[1] = 100;
- round_factors[2] = 10;
- round_factors[3] = 7; // Beyond this point, the error function is linear
- round_factors[4] = 3;
- round_factors[5] = 1;
- round_factors[6] = 0.7;
- round_factors[7] = 0.3;
- round_factors[8] = 0.1;
- round_factors[9] = 0.07;
- round_factors[10] = 0.03;
- round_factors[11] = 0.01;
-
- // THINK: Considering using error magnitudes in this scenario
-
-
- float round_factor = round_factors[error_scale];
- INFO("round_factor = %f \n", round_factor);
-
- hostToDeviceCopy(x);
-
- int blockSize = 128;
- int gridSize = (int) ceil ((float) x->num_elems / blockSize);
- INFO("blockSize = %d, gridSize = %d \n", blockSize, gridSize);
-
- // NOTE: Check if a large gridSize will work with really large tensors
- vecConstMul<<<gridSize, blockSize>>>((float*) x->gpu_data, round_factor, x->num_elems);
- //vecRound<<<gridSize, blockSize>>>((float*) x->gpu_data, x->num_elems);
-
- deviceToHostCopy(x);
- randomCeilAndFloor((float*) x->host_data, x->num_elems);
- hostToDeviceCopy(x);
-
- vecConstDiv<<<gridSize, blockSize>>>((float*) x->gpu_data, round_factor, x->num_elems);
-
- Norm_t* norms = calculateNorms2(x, x_original);
-
- profileEvent("tensorRoundError_end", true);
-
- return (void*) norms;
-}
-
-
-
-
-// Routine for Adding Gaussian Error
-void* addGaussianError(void* x_ptr, int error_scale){
-
- if(error_scale > 20 || error_scale < 0){
- ERROR("Error Scale out of bounds \n");
- }
-
- INFO("*** TensorAddError \n");
- profileEvent("tensorAddError");
-
- Tensor* x = (Tensor*) x_ptr;
-
- size_t* dim_sizes = x->dims.dim_sizes;
- Tensor* bias = (Tensor*) create4DTensor(x->cur_type, x->data_format,
- dim_sizes[0], dim_sizes[1],
- dim_sizes[2], dim_sizes[3]);
-
- Tensor* x_original = (Tensor*) create4DTensor(x->cur_type, x->data_format,
- dim_sizes[0], dim_sizes[1],
- dim_sizes[2], dim_sizes[3]);
-
- // Copying x data into x_original - for computing Norms
- tensorCopy(x, x_original);
-
- // NOTE: Error scale is used to generate the bias matrix
- initRandValues(bias, error_scale);
-
- hostToDeviceCopy(x);
- //hostToDeviceCopy(bias);
-
-
- int blockSize = 1024;
- int gridSize = (int) ceil ((float) x->num_elems / blockSize);
- INFO("blockSize = %d, gridSize = %d \n", blockSize, gridSize);
-
- // NOTE: Check if a large gridSize will work with really large tensors
- vecMul<<<gridSize, blockSize>>>((float*) x->gpu_data, (float*) bias->gpu_data, x->num_elems);
-
- float alpha = 1.0f;
-
- // FIXIT: routine fails for 3D tensors
- checkCUDNN(cudnnAddTensor(cudnnHandle, &alpha, bias->tensor_desc,
- bias->gpu_data, &alpha, x->tensor_desc, x->gpu_data));
-
-
- //Norm_t* norms = calculateNorms2(x, x_original);
- //Norm_t* norms = calculateNormsGPU(x, x_original);
-
- Norm_t* norms = calculateNormsTreeReduction(x, x_original);
-
- freeTensor(x_original);
- freeTensor(bias);
-
-
- profileEvent("tensorAddError_end", true);
-
- return (void*) norms;
-}
-
-
-
void initPromiseRandValues(Tensor* bias, int error_scale){
float scaling_values[10];
@@ -983,18 +682,6 @@ void* quantizeTensorPromise(void* input_ptr, float min, float max){
return input;
}
-
-void* tensorAddError(void* x_ptr, int error_scale){
-
- void * new_x = addGaussianError(x_ptr, error_scale);
- //void * new_x = addRoundError(x_ptr, error_scale);
- //void * new_x = addBitError(x_ptr, error_scale);
- return new_x;
-}
-
-
-
-
}
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/group_conv.cu b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/group_conv.cu
index 4b49a3702b1938ceed9829cc3572474c7cb82420..ab8896369ab4a8b4b6e2bc44a3896034001300cc 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/group_conv.cu
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/group_conv.cu
@@ -397,28 +397,6 @@ void* tensorConvCutlass(void* input_ptr, void* filter_ptr,
cudaDeviceSynchronize();
profileEvent("Conv_end", true);
-
- #ifdef ERROR_INJECTION_ENABLED
-
- if (op_counter >= total_ops) {
- ERROR("No accuracy flag found \n");
- }
-
- int op_acc = op_accuracies[op_counter];
-
- // Skip errorInjection if explicitly requested
- if (skip_tensors.find(op_counter) != skip_tensors.end()) {
- op_acc = 0;
- }
-
- void* error_norms = tensorAddError(output, op_acc);
- add_norms(error_norms, "tensorConv", op_acc);
- add_conv_overheads(input, filter, vertical_stride, horizontal_stride, op_acc);
-
- op_counter++;
-
- #endif
-
return output;
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/init_api.cc b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/init_api.cc
index 8b5c4aaf93db40c038c4a9a30569318ae00d6be1..74ee15c2dcb916f9a7a24fdc1318255e626844b3 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/init_api.cc
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/init_api.cc
@@ -50,10 +50,6 @@ void llvm_hpvm_initTensorRt(int gpuid) {
#endif
-#ifdef ERROR_INJECTION_ENABLED
- readOpenTunerFlags("opentuner_flags");
-#endif
-
runtime_initialized = true;
}
@@ -73,11 +69,6 @@ void llvm_hpvm_initApproxhpvmRt(int gpuid) {
void llvm_hpvm_cleanupApproxhpvmRt() {}
void dumpAccuracyNorms() {
-
-#ifdef ERROR_INJECTION_ENABLED
-
-#endif
-
dump_result("accuracy_summary");
}
diff --git a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_runtime.cu b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_runtime.cu
index 319936b482c455af2fcc0280adb15d7c126c088a..1b28ccaa191ba3e52255b934c894e543dd052773 100644
--- a/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_runtime.cu
+++ b/hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_runtime.cu
@@ -81,21 +81,6 @@ void* tensorAdd(void* x_ptr, void* bias_ptr){
profileEvent("Add_end", true);
- #ifdef ERROR_INJECTION_ENABLED
- if(op_counter >= total_ops){
- ERROR("No accuracy flag found \n");
- }
-
- int op_acc = op_accuracies[op_counter];
-
- void* error_norms = tensorAddError(x, op_acc);
- add_norms(error_norms, "tensorAdd", op_acc);
- add_bias_overheads(x, op_acc);
- op_counter++;
-
- #endif
-
-
return x;
}
@@ -241,25 +226,6 @@ void* tensorConvolution(void* input_ptr, void* filter_ptr,
&beta, output->tensor_desc, output->gpu_data));
profileEvent("Conv_end", true);
-
-
- #ifdef ERROR_INJECTION_ENABLED
-
- if(op_counter >= total_ops){
- ERROR("No accuracy flag found \n");
- }
-
- int op_acc = op_accuracies[op_counter];
-
- void* error_norms = tensorAddError(output, op_acc);
- add_norms(error_norms, "tensorConv", op_acc);
- add_conv_overheads(input, filter, vertical_stride, horizontal_stride, op_acc);
-
- op_counter++;
-
- #endif
-
-
return output;
}
@@ -328,24 +294,6 @@ void* tensorPooling(void* input_ptr,
input->gpu_data, &beta, output->tensor_desc, output->gpu_data));
profileEvent("Pool_end", true);
-
-
- #ifdef ERROR_INJECTION_ENABLED
-
- if(op_counter >= total_ops){
- ERROR("No accuracy flag found \n");
- }
-
- int op_acc = op_accuracies[op_counter];
- void* error_norms = tensorAddError(output, op_acc);
- add_norms(error_norms, "tensorPooling", op_acc);
- add_pool_overheads(input, window_height, vertical_stride, op_acc);
-
- op_counter++;
-
- #endif
-
-
return output;
}
@@ -420,26 +368,6 @@ void* tensorGemmGPU(void* lhs_ptr, void* rhs_ptr ){
profileEvent("Mul_end", true);
-
-
-
- #ifdef ERROR_INJECTION_ENABLED
-
- if(op_counter >= total_ops){
- ERROR("No accuracy flag found \n");
- }
-
- int op_acc = op_accuracies[op_counter];
-
- void* error_norms = tensorAddError(output, op_acc);
- add_norms(error_norms, "tensorGemm", op_acc);
- add_gemm_overheads(lhs_ptr, rhs_ptr, op_acc);
-
- op_counter++;
-
- #endif
-
-
return output;
}
@@ -473,23 +401,6 @@ void* tensorRelu(void* input_ptr){
input->tensor_desc, input->gpu_data));
profileEvent("Relu_end", true);
-
-
- #ifdef ERROR_INJECTION_ENABLED
-
- if(op_counter >= total_ops){
- ERROR("No accuracy flag found \n");
- }
-
- int op_acc = op_accuracies[op_counter];
-
- void* error_norms = tensorAddError(input, op_acc);
- add_norms(error_norms, "tensorRelu", op_acc);
- add_relu_overheads(input, op_acc);
- op_counter++;
- #endif
-
-
return input;
}
@@ -546,22 +457,6 @@ void* tensorRelu2(void* input_ptr, float min, float max){
profileEvent("Relu_end", true);
-
-
- #ifdef ERROR_INJECTION_ENABLED
-
- if(op_counter >= total_ops){
- ERROR("No accuracy flag found \n");
- }
-
- int op_acc = op_accuracies[op_counter];
- void* error_norms = tensorAddError(input, op_acc);
- add_norms(error_norms, "tensorClippedRelu", op_acc);
- add_relu_overheads(input, op_acc);
- op_counter++;
- #endif
-
-
return input;
}
@@ -591,22 +486,6 @@ void* tensorTanh(void* input_ptr){
input->tensor_desc, input->gpu_data));
profileEvent("Tanh_end", true);
-
-
- #ifdef ERROR_INJECTION_ENABLED
-
- if(op_counter >= total_ops){
- ERROR("No accuracy flag found \n");
- }
-
- int op_acc = op_accuracies[op_counter];
- void* error_norms = tensorAddError(input, op_acc);
- add_norms(error_norms, "tensorTanh", op_acc);
- add_relu_overheads(input, op_acc);
- op_counter++;
- #endif
-
-
return input;
}
@@ -650,22 +529,6 @@ void* tensorBatchNorm(void* input_ptr, void* gamma_ptr, void* beta_ptr,
epsilon));
profileEvent("BatchNorm_end", true);
-
-
- #ifdef ERROR_INJECTION_ENABLED
-
- if(op_counter >= total_ops){
- ERROR("No accuracy flag found \n");
- }
-
- int op_acc = op_accuracies[op_counter];
- void* error_norms = tensorAddError(input, op_acc);
- add_norms(error_norms, "tensorBatchNorm", op_acc);
- add_relu_overheads(input, op_acc);
- op_counter++;
- #endif
-
-
return input;
}