diff --git a/common/arch.h b/common/arch.h index 13a1983..012c28e 100644 --- a/common/arch.h +++ b/common/arch.h @@ -211,7 +211,13 @@ enum InstructionType { // List of instruction groups enum GroupType { G_FP64 = 0, // FP64 arithmetic instructions - G_FP32, // FP32 arithmetic instructions + G_FP32, // FP32 arithmetic instructions + /** + * Fernando Fernandes, 10/2022 + * Add the FP16 from FP16 MMA fault sites separated from FP32 + */ + G_FP16, + G_MMA, G_LD, // instructions that read from emory G_PR, // instructions that write to PR registers only G_NODEST, // instructions with no destination register @@ -229,7 +235,13 @@ enum BitFlipModel { FLIP_TWO_BITS, // flip two adjacent bits RANDOM_VALUE, // write a random value. ZERO_VALUE, // write value 0 - NUM_BFM_TYPES + /** + * Fernando Fernandes, 10/2022 + * Add warp wide fault models + */ + WARP_RANDOM_VALUE, // random in a warp + WARP_ZERO_VALUE, // zero in all the warp + NUM_BFM_TYPES }; diff --git a/common/globals.h b/common/globals.h index 2853a53..3cc8d74 100644 --- a/common/globals.h +++ b/common/globals.h @@ -110,8 +110,12 @@ const char * instTypeNames[NUM_ISA_INSTRUCTIONS] = { "GETLMEMBASE", "SETCRSPTR", "SETLMEMBASE" , "PMTRIG", "SETCTAID" }; +/** +* Fernando Fernandes, 10/2022 +* Add the FP16 and MMA fault sites separated from FP32 +*/ const char * instGrouptNames[NUM_INST_GROUPS] = { - "fp64", "fp32", "ld", "pr", "nodest", "others", "gppr", "gp" + "fp64", "fp32", "fp16", "mma", "ld", "pr", "nodest", "others", "gppr", "gp" }; int fp64Inst[] = { @@ -122,6 +126,18 @@ int fp32Inst[] = { FADD, FADD32I, FCMP, FFMA, FFMA32I, FMNMX, FMUL, FMUL32I, FSEL, FSET, FSWZADD, IPA, DSET }; +/** +* Fernando Fernandes, 09/2022 +* Add the FP16 and MMA fault sites separated from FP32 +*/ +int fp16Inst[] = { + HADD2, HADD2_32I, HFMA2, HFMA2_32I, HMUL2, HMUL2_32I, HSET2, HSETP2, +}; + +int MMAInst[] = { + IMMA, HMMA, +}; + int ldInst[] = { LD, LDC, LDG, LDL, LDS, SULD, SUST, TLD, TLD4, TLD4S, TLDS }; @@ -142,11 +158,15 @@ int noDestInst[] = { int otherInst[] = { // Floating-point Instructions - MUFU, RRO, HADD2, HADD2_32I, HFMA2, HFMA2_32I, HMUL2, HMUL2_32I, HSET2, HSETP2, +/** +* Fernando Fernandes, 09/2022 +* Add the FP16 and MMA fault sites separated from FP32 +*/ + MUFU, RRO, // HADD2, HADD2_32I, HFMA2, HFMA2_32I, HMUL2, HMUL2_32I, HSET2, HSETP2, // Integer Instructions IDP, IDP4A, BFE, BFI, BMSK, BREV, FLO, IADD, IADD3, IADD32I, ICMP, IMAD, IMAD32I, IMADSP, IMNMX, IMUL, IMUL32I, ISCADD, ISCADD32I, ISET, LEA, LOP, LOP3, LOP32I, POPC, SHF, SHL, SHR, XMAD, // MMA instructions - IMMA, HMMA, +// IMMA, HMMA, // Video Instructions VABSDIFF, VADD, VMAD, VMNMX, VSET, VSHL, VSHR, VABSDIFF4, // Conversion Instructions @@ -249,6 +269,14 @@ int getOpGroupNum(int opcode) { return G_FP64; if (checkOpType(opcode, fp32Inst, sizeof(fp32Inst)/sizeof(int))) return G_FP32; + /** + * Fernando Fernandes, 09/2022 + * Add the FP16 and MMA fault sites separated from FP32 + */ + if (checkOpType(opcode, fp16Inst, sizeof(fp16Inst) / sizeof(int))) + return G_FP16; + if (checkOpType(opcode, MMAInst, sizeof(MMAInst) / sizeof(int))) + return G_MMA; if (checkOpType(opcode, ldInst, sizeof(ldInst)/sizeof(int))) return G_LD; if (checkOpType(opcode, prOnlyInst, sizeof(prOnlyInst)/sizeof(int))) diff --git a/injector/Makefile b/injector/Makefile index 5edba0b..3db1f87 100644 --- a/injector/Makefile +++ b/injector/Makefile @@ -17,7 +17,7 @@ DUMMY=0 SOURCES=$(wildcard *.cu) OBJECTS=$(SOURCES:.cu=.o) -ARCH=35 +ARCH=70 NVBIT_TOOL=injector.so diff --git a/injector/inject_funcs.cu b/injector/inject_funcs.cu index ded91b6..8afc22d 100644 --- a/injector/inject_funcs.cu +++ b/injector/inject_funcs.cu @@ -34,19 +34,30 @@ __inline__ __device__ int get_flat_tid() { // Get bit-mask for error injection. Old value will be XORed with this mask later to inject the error. __inline__ __device__ unsigned int get_mask(uint32_t bitFlipModel, float bitIDSeed, unsigned int oldVal) { + /** + * Fernando Fernandes, 10/2022 + * Added the mask creation for warp wide injections + */ return (bitFlipModel == FLIP_SINGLE_BIT) * ((unsigned int)1<<(int)(32*bitIDSeed)) + (bitFlipModel == FLIP_TWO_BITS) * ((unsigned int)3<<(int)(31*bitIDSeed)) + - (bitFlipModel == RANDOM_VALUE) * (((unsigned int)-1) * bitIDSeed) + - (bitFlipModel == ZERO_VALUE) * oldVal; + (bitFlipModel == RANDOM_VALUE || bitFlipModel == WARP_RANDOM_VALUE) * (((unsigned int)-1) * bitIDSeed) + + (bitFlipModel == ZERO_VALUE || bitFlipModel == WARP_ZERO_VALUE) * oldVal; } extern "C" __device__ __noinline__ void inject_error(uint64_t piinfo, uint64_t pcounters, uint64_t pverbose_device, int offset, int index, int grp_index, int predicate, int destGPRNum, int regval, int numDestGPRs, int destPRNum1, int destPRNum2, int maxRegs) { - inj_info_t* inj_info = (inj_info_t*)piinfo; + inj_info_t* inj_info = (inj_info_t*)piinfo; uint32_t verbose_device = *((uint32_t *)pverbose_device); uint64_t * counters = (uint64_t *)pcounters; + // Fernando Fernandes, 10/2022: The inj_info is now an array of 32 position of inj_info_t + // each position for each lane_id. This is only valid for warp wide injections, other fault models use position 0 + const unsigned lane_id = get_laneid(); + const bool warp_wide_fi = inj_info->bitFlipModel == WARP_ZERO_VALUE || inj_info->bitFlipModel == WARP_RANDOM_VALUE; + if (warp_wide_fi) { + inj_info += lane_id; + } if (verbose_device) inj_info->debug[NUM_DEBUG_VALS-1] = 1; @@ -75,6 +86,12 @@ extern "C" __device__ __noinline__ void inject_error(uint64_t piinfo, uint64_t p bool injectFlag = false; switch (igid) { + /** + * Fernando Fernandes, 10/2022 + * Enable the FP16 and the MMA injections + */ + case G_FP16: // It is supposed to work for FP16 and FP16MMA, as they have destination registers + case G_MMA: case G_FP32: // inject into one of the dest reg case G_FP64: // inject into one of the regs written by the inst case G_LD: // inject into one of the regs written by the inst @@ -96,14 +113,17 @@ extern "C" __device__ __noinline__ void inject_error(uint64_t piinfo, uint64_t p case G_NODEST: // do nothing default: break; } - + // Fernando Fernandes, 10/2022: + // If the warp injections is selected, sync between all the active threads within the warp + if (warp_wide_fi) + injectFlag = __any_sync(__activemask(), injectFlag) != 0; if (verbose_device && injectFlag) printf("inj_info->instID=%ld, %ld, %ld, %ld\n", inj_info->instID, currCounter1, currCounter2, currCounter3); if (injectFlag) { // assert(0 == 10); if (verbose_device) - printf("offset=0x%x, igid:%d, destGPRNum=%d, grp_index=%d\n", offset, igid, destGPRNum, grp_index); + printf("offset=0x%x, igid:%d, destGPRNum=%d, grp_index=%d laneId=%d\n", offset, igid, destGPRNum, grp_index, lane_id); // Fernando Fernandes, 10/2022: debug message // We need to randomly select one register from numDestGPRs + (destPRNum1 != -1) + (destPRNum2 != -1) int totalDest = numDestGPRs + (destPRNum1 != -1) + (destPRNum2 != -1); assert(totalDest > 0); @@ -142,7 +162,7 @@ extern "C" __device__ __noinline__ void inject_error(uint64_t piinfo, uint64_t p assert(inj_info->debug[12] == inj_info->opcode); assert(inj_info->debug[13] == inj_info->pcOffset); if (verbose_device) - printf("done here\n"); + printf("done here, lane:%d\n", lane_id); // Fernando Fernandes, 10/2022: debug message } else { assert(0 == 2); } diff --git a/injector/injector.cu b/injector/injector.cu index 9dfd65e..8d5cfb6 100644 --- a/injector/injector.cu +++ b/injector/injector.cu @@ -46,9 +46,15 @@ std::string injInputFilename = "nvbitfi-injection-info.txt"; pthread_mutex_t mutex; -__managed__ inj_info_t inj_info; +/** +* Fernando Fernandes, 10/2022 +* Add warp wide fault models +*/ +#define WARP_SIZE 32 +__managed__ inj_info_t inj_info_warp[WARP_SIZE]; void reset_inj_info() { + inj_info_t inj_info; inj_info.areParamsReady = false; inj_info.kernelName[0] = '\0'; inj_info.kernelCount = -1; @@ -68,9 +74,10 @@ void reset_inj_info() { for (int i=0; i &instrs = nvbit_get_instrs(ctx, f); if (verbose) - printf(":::NVBit-inject-error; NOT inspecting: %s; %d, %d, num_static_instrs: %ld; maxregs: %d:::", kname.c_str(), kernel_id, inj_info.kernelCount, instrs.size(), get_maxregs(f)); + printf(":::NVBit-inject-error; NOT inspecting: %s; %d, %d, num_static_instrs: %ld; maxregs: %d:::", kname.c_str(), kernel_id, inj_info_warp[0].kernelCount, instrs.size(), get_maxregs(f)); } } } @@ -353,13 +366,17 @@ void nvbit_at_cuda_event(CUcontext ctx, int is_exit, nvbit_api_cuda_t cbid, parse_params(injInputFilename); // injParams are updated based on injection seed file // print_inj_info(); - inj_info.errorInjected = false; - inj_info.areParamsReady = (inj_info.kernelCount== kernel_id); // areParamsReady = true for the selected kernel - if (verbose) inj_info.debug[NUM_DEBUG_VALS-1] = -1; // set debug flag to check whether the the instrumented kernel was executed - if (verbose) printf("setting areParamsReady=%d, inj_info.kernelCount=%d, kernel_id=%d\n", inj_info.areParamsReady, inj_info.kernelCount, kernel_id); + // Set all the array of inj_info with the same info + for(auto lane_id = 0; lane_id < WARP_SIZE; lane_id++) { + auto& inj_info = inj_info_warp[lane_id]; + inj_info.errorInjected = false; + inj_info.areParamsReady = (inj_info.kernelCount== kernel_id); // areParamsReady = true for the selected kernel + if (verbose) inj_info.debug[NUM_DEBUG_VALS-1] = -1; // set debug flag to check whether the the instrumented kernel was executed + if (verbose) printf("setting areParamsReady=%d, inj_info.kernelCount=%d, kernel_id=%d, laneId=%d\n", inj_info.areParamsReady, inj_info.kernelCount, kernel_id, lane_id); + } cudaDeviceSynchronize(); - nvbit_enable_instrumented(ctx, p->f, inj_info.areParamsReady); // should we run the un-instrumented code? + nvbit_enable_instrumented(ctx, p->f, inj_info_warp[0].areParamsReady); // should we run the un-instrumented code? // nvbit_enable_instrumented(ctx, p->f, false); // for debugging cudaDeviceSynchronize(); } else { @@ -376,48 +393,57 @@ void nvbit_at_cuda_event(CUcontext ctx, int is_exit, nvbit_api_cuda_t cbid, } std::string kname = removeSpaces(nvbit_get_func_name(ctx,p->f)); - if (inj_info.areParamsReady) { - inj_info.areParamsReady = false; - int num_ctas = 0; - if ( cbid == API_CUDA_cuLaunchKernel_ptsz || cbid == API_CUDA_cuLaunchKernel) { - cuLaunchKernel_params * p2 = (cuLaunchKernel_params*) params; - num_ctas = p2->gridDimX * p2->gridDimY * p2->gridDimZ; - } - assert(fout.good()); - fout << "Injection data" << std::endl; - fout << "index: " << kernel_id << std::endl; - fout << "kernel_name: " << kname << std::endl; - fout << "ctas: " << num_ctas << std::endl; - fout << "instrs: " << get_inst_count() << std::endl; - - write_inj_info(); - - if (inj_info.opcode == NOP) { - fout << "Error not injected\n"; - } - - if (verbose != 0 && inj_info.debug[2] != inj_info.debug[3]) { // sanity check - fout << "ERROR FAIL in kernel execution; Expected reg value doesn't match; \n"; - fout << "maxRegs: " << inj_info.debug[0] << ", destGPRNum: " << inj_info.debug[1] << ", expected_val: " - << std::hex << inj_info.debug[2] << ", myval: " << inj_info.debug[3] << std::dec << "\n"; - fout << std::endl; - std::cout << "NVBit-inject-error; ERROR FAIL in kernel execution; Expected reg value doesn't match; \n"; - std::cout << "maxRegs: " << inj_info.debug[0] << ", destGPRNum: " << inj_info.debug[1] << ", expected_val: " - << std::hex << inj_info.debug[2] << ", myval: " << inj_info.debug[3] << std::dec << "\n"; - for (int x=4; x<10; x++) { - std::cout << "debug[" << x << "]: " << std::hex << inj_info.debug[x] << "\n"; - } - std::cout << "debug[11]: " << std::hex << inj_info.debug[11] << "\n"; - std::cout << "debug[12]: " << inj_info.debug[12] << " " << instTypeNames[inj_info.debug[12]]<< "\n"; - std::cout << "debug[13]: " << inj_info.debug[13] << "\n"; - std::cout << "debug[14]: " << std::hex << inj_info.debug[14] << "\n"; - assert(inj_info.debug[2] == inj_info.debug[3]); - // printf("\nmaxRegs: %d, destGPRNum: %d, expected_val: %x, myval: %x, myval@-1: %x, myval@+1: %x, myval with maxRegs+1: %x, myval with maxRegs-1: %x\n", - // inj_info.debug[0], inj_info.debug[1], inj_info.debug[2], inj_info.debug[3], inj_info.debug[4], inj_info.debug[5], inj_info.debug[6], inj_info.debug[7]); - } - fout.flush(); - } - if (verbose) printf("\n index: %d; kernel_name: %s; used_instrumented=%d; \n", kernel_id, kname.c_str(), inj_info.debug[NUM_DEBUG_VALS-1]); + /** + * Fernando Fernandes, 10/2022 + * if a warp fault model is selected verify all the elements in the inj_info_warp array + */ + for(auto lane_id = 0; lane_id < WARP_SIZE; lane_id++) { + auto& inj_info = inj_info_warp[lane_id]; + if (inj_info.areParamsReady) { + inj_info.areParamsReady = false; + int num_ctas = 0; + if ( cbid == API_CUDA_cuLaunchKernel_ptsz || cbid == API_CUDA_cuLaunchKernel) { + cuLaunchKernel_params * p2 = (cuLaunchKernel_params*) params; + num_ctas = p2->gridDimX * p2->gridDimY * p2->gridDimZ; + } + assert(fout.good()); + fout << "Injection data, laneId:" << lane_id << std::endl; + fout << "index: " << kernel_id << std::endl; + fout << "kernel_name: " << kname << std::endl; + fout << "ctas: " << num_ctas << std::endl; + fout << "instrs: " << get_inst_count() << std::endl; + + write_inj_info(inj_info); + + if (inj_info.opcode == NOP) { + fout << "Error not injected\n"; + } + + if (verbose != 0 && inj_info.debug[2] != inj_info.debug[3]) { // sanity check + fout << "ERROR FAIL in kernel execution; Expected reg value doesn't match; \n"; + fout << "maxRegs: " << inj_info.debug[0] << ", destGPRNum: " << inj_info.debug[1] << ", expected_val: " + << std::hex << inj_info.debug[2] << ", myval: " << inj_info.debug[3] << std::dec << "\n"; + fout << std::endl; + std::cout << "NVBit-inject-error; ERROR FAIL in kernel execution; Expected reg value doesn't match; \n"; + std::cout << "maxRegs: " << inj_info.debug[0] << ", destGPRNum: " << inj_info.debug[1] << ", expected_val: " + << std::hex << inj_info.debug[2] << ", myval: " << inj_info.debug[3] << std::dec << "\n"; + for (int x=4; x<10; x++) { + std::cout << "debug[" << x << "]: " << std::hex << inj_info.debug[x] << "\n"; + } + std::cout << "debug[11]: " << std::hex << inj_info.debug[11] << "\n"; + std::cout << "debug[12]: " << inj_info.debug[12] << " " << instTypeNames[inj_info.debug[12]]<< "\n"; + std::cout << "debug[13]: " << inj_info.debug[13] << "\n"; + std::cout << "debug[14]: " << std::hex << inj_info.debug[14] << "\n"; + assert(inj_info.debug[2] == inj_info.debug[3]); + // printf("\nmaxRegs: %d, destGPRNum: %d, expected_val: %x, myval: %x, myval@-1: %x, myval@+1: %x, myval with maxRegs+1: %x, myval with maxRegs-1: %x\n", + // inj_info.debug[0], inj_info.debug[1], inj_info.debug[2], inj_info.debug[3], inj_info.debug[4], inj_info.debug[5], inj_info.debug[6], inj_info.debug[7]); + } + fout.flush(); + } + if (verbose) printf("\n index: %d; kernel_name: %s; used_instrumented=%d; \n", kernel_id, kname.c_str(), inj_info.debug[NUM_DEBUG_VALS-1]); + if (inj_info.bitFlipModel != WARP_ZERO_VALUE && inj_info.bitFlipModel != WARP_RANDOM_VALUE) + break; + } kernel_id++; // always increment kernel_id on kernel exit cudaDeviceSynchronize(); diff --git a/profiler/Makefile b/profiler/Makefile index a542a7f..779c271 100644 --- a/profiler/Makefile +++ b/profiler/Makefile @@ -18,7 +18,7 @@ NVCC_PATH=-L $(subst bin/nvcc,lib64,$(shell which nvcc | tr -s /)) SOURCES=$(wildcard *.cu) OBJECTS=$(SOURCES:.cu=.o) -ARCH=35 +ARCH=70 # Profiling is often slow. Using an approximate profile can speed it up by # orders of magnitude. Set FAST_APPROXIMATE_PROFILE to SKIP_PROFILED_KERNELS diff --git a/profiler/inject_funcs.cu b/profiler/inject_funcs.cu index 5778300..5f2eecb 100644 --- a/profiler/inject_funcs.cu +++ b/profiler/inject_funcs.cu @@ -19,6 +19,17 @@ #include "utils/utils.h" #include "arch.h" +/** +* Fernando Fernandes, 10/2022 +* Add the support for ballot_sync for architectures >= 70 +*/ +__device__ inline int profiler_ballot(int pred) { +#if __CUDA_ARCH__ >= 700 + return __ballot_sync(0xFFFFFFFF, pred); +#else + return ballot(pred); +#endif +} // Global counters are incremented once per warp extern "C" __device__ __noinline__ void count_instrs(uint64_t pcounters, int index, int grp_index, int predicate, int num_counters) { @@ -27,12 +38,12 @@ extern "C" __device__ __noinline__ void count_instrs(uint64_t pcounters, int ind // Optimization: Instead of all the threads in a warp performing atomicAdd, // let's count the number of active threads with predicate=1 in a warp and let just one thread // (leader) in the warp perform the atomicAdd - const int active_mask = ballot(1); + const int active_mask = profiler_ballot(1); const int leader = __ffs(active_mask) - 1; const int laneid = get_laneid(); // compute the predicate mask - const int predicate_mask = ballot(predicate); + const int predicate_mask = profiler_ballot(1); //ballot(predicate); const int num_threads = __popc(predicate_mask); if (laneid == leader) { // Am I the leader thread diff --git a/scripts/params.py b/scripts/params.py index 7aa88b7..59b70e1 100644 --- a/scripts/params.py +++ b/scripts/params.py @@ -96,15 +96,18 @@ ####################################################################### G_FP64 = 0 G_FP32 = 1 -G_LD = 2 -G_PR = 3 -G_NODEST = 4 # not really an igid -G_OTHERS = 5 -G_GPPR = 6 # instructions that write to either a GPR or a PR register -G_GP = 7 # instructions that write to a GPR register -NUM_INST_GROUPS = 8 +G_FP16 = 2 # Fernando Fernandes, 10/2022 New fault sites +G_MMA = 3 +G_LD = 4 +G_PR = 5 +G_NODEST = 6 # not really an igid +G_OTHERS = 7 +G_GPPR = 8 # instructions that write to either a GPR or a PR register +G_GP = 9 # instructions that write to a GPR register +NUM_INST_GROUPS = 10 -IGID_STR = [ "fp64", "fp32", "ld", "pr", "nodest", "others", "gppr", "gp" ] +# Fernando Fernandes, 10/2022 New fault sites +IGID_STR = [ "fp64", "fp32", "fp16", "mma", "ld", "pr", "nodest", "others", "gppr", "gp" ] ####################################################################### @@ -115,8 +118,11 @@ FLIP_TWO_BITS = 1 RANDOM_VALUE = 2 ZERO_VALUE = 3 - -EM_STR = [ "FLIP_SINGLE_BIT", "FLIP_TWO_BITS", "RANDOM_VALUE", "ZERO_VALUE"] +# Fernando Fernandes, 10/2022 +# Add warp wide fault models +WARP_RANDOM_VALUE = 4 +WARP_ZERO_VALUE = 5 +EM_STR = [ "FLIP_SINGLE_BIT", "FLIP_TWO_BITS", "RANDOM_VALUE", "ZERO_VALUE", "WARP_RANDOM_VALUE", "WARP_ZERO_VALUE"] rf_inst = "" @@ -180,11 +186,12 @@ G_GP: [FLIP_SINGLE_BIT] # Supported models -# G_GP: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE] -# G_FP64: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE] -# G_FP32: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE] -# G_LD: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE] - +# G_GP: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE, WARP_RANDOM_VALUE, WARP_ZERO_VALUE], + # G_FP64: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE, WARP_RANDOM_VALUE, WARP_ZERO_VALUE], + # G_FP32: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE, WARP_RANDOM_VALUE, WARP_ZERO_VALUE], + # G_LD: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE, WARP_RANDOM_VALUE, WARP_ZERO_VALUE], + # G_FP16: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE, WARP_RANDOM_VALUE, WARP_ZERO_VALUE], + # G_MMA: [FLIP_SINGLE_BIT, FLIP_TWO_BITS, RANDOM_VALUE, ZERO_VALUE, WARP_RANDOM_VALUE, WARP_ZERO_VALUE], } @@ -208,6 +215,15 @@ 1, # expected runtime "" # additional parameters to the run.sh ], + # Fernando Fernandes, 10/2022 New fault sites + # A simple example with CUBLAS to test the FP16 and MMA injections + 'simple_cublas': [ + NVBITFI_HOME + '/test-apps/simple_cublas', # workload directory + 'simple_cublas', # binary name + NVBITFI_HOME + '/test-apps/simple_cublas/', # path to the binary file + 1, # expected runtime + "" # additional parameters to the run.sh + ], } ######################################################################### diff --git a/test-apps/simple_cublas/Makefile b/test-apps/simple_cublas/Makefile new file mode 100644 index 0000000..cb6ca7a --- /dev/null +++ b/test-apps/simple_cublas/Makefile @@ -0,0 +1,33 @@ +TARGET = simple_cublas +#USE_TENSOR_CORES=0 +#USE_TENSOR_CORES=1 + +# CUDA +CUDA_PATH=/usr/local/cuda +NVCC = $(CUDA_PATH)/bin/nvcc + +SMS = 70 +$(foreach sm,$(SMS), $(eval ARCH += -gencode arch=compute_$(sm),code=sm_$(sm))) + +FLAGS = -Xptxas -v $(ARCH) -L$(CUDA_PATH)/lib64 -lcublas + +ifeq ($(USE_TENSOR_CORES), 1) +FLAGS+= -DUSE_TENSOR_CORES=1 +endif + +all: clean $(TARGET) + +$(TARGET): + $(NVCC) -o $(TARGET) $(FLAGS) $(TARGET).cu + +test: + ./$(TARGET) + +golden: + ./$(TARGET) >golden_stdout.txt 2>golden_stderr.txt + +clean: + rm -f *.o *~ $(TARGET) + +clobber: clean + rm -f golden* *_regcount.p stdout* stderr* diff --git a/test-apps/simple_cublas/run.sh b/test-apps/simple_cublas/run.sh new file mode 100755 index 0000000..a177207 --- /dev/null +++ b/test-apps/simple_cublas/run.sh @@ -0,0 +1,2 @@ +#!/bin/bash +eval ${PRELOAD_FLAG} ${BIN_DIR}/simple_cublas > stdout.txt 2> stderr.txt diff --git a/test-apps/simple_cublas/sdc_check.sh b/test-apps/simple_cublas/sdc_check.sh new file mode 100755 index 0000000..9d97eb6 --- /dev/null +++ b/test-apps/simple_cublas/sdc_check.sh @@ -0,0 +1,51 @@ +#!/bin/bash +# +# Copyright 2020, NVIDIA CORPORATION. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +# + +################################################################################## +# Brief explanation of the log files: +# (1) diff.log: If the program generates a separate output file (e.g., +# output.txt), diff the file with the golden output file and store the diff in +# diff.log. The program output should be deterministic. Please exclude +# non-deterministic values (e.g., runtimes) from the file before you use +# `diff` or similar utilities. +# (2) stdout_diff.log: diff stdout generated by the program with # the golden +# stdout and store it in stdout_diff.log. Remove the string that matches +# ":::Injecting.*:::" from the stdout. +# (3) stderr_diff.log: diff stderr generated by the program with the golden +# stderr and store it in stderr_diff.log. +# (4) special_check.log: if the application specific check fails, the +# special_check.log should contain some non-empty string. +# +# If the user prefers to skip one of these checks, he/she should create an +# empty log file (e.g., touch diff.log). +################################################################################## + +# if your program creates an output file (e.g., output.txt) compare it to the file created just now and store the difference in diff.log +# Example: diff output.txt ${APP_DIR}/golden_output.txt > diff.log +touch diff.log + +# comparing stdout generated by your program +diff stdout.txt ${APP_DIR}/golden_stdout.txt > stdout_diff.log + +# comparing stderr generated by your program +diff stderr.txt ${APP_DIR}/golden_stderr.txt > stderr_diff.log + +# Application specific output: The following check will be performed only if at least one of diff.log, stdout_diff.log, and stderr_diff.log is different +grep sum stdout.txt > selected_output.txt +grep sum ${APP_DIR}/golden_stdout.txt > selected_golden_output.txt +diff selected_output.txt selected_golden_output.txt > special_check.log + diff --git a/test-apps/simple_cublas/simple_cublas.cu b/test-apps/simple_cublas/simple_cublas.cu new file mode 100644 index 0000000..7ddff9d --- /dev/null +++ b/test-apps/simple_cublas/simple_cublas.cu @@ -0,0 +1,81 @@ +#include +#include +#include +#include + +#if USE_TENSOR_CORES == 1 +#define BLAS_MODE CUBLAS_TF32_TENSOR_OP_MATH +#else +#define BLAS_MODE CUBLAS_PEDANTIC_MATH +#endif + +static void checkCublasErrorsCall(cublasStatus_t error, int line, const char *file) { + if (error == CUBLAS_STATUS_SUCCESS) { + return; + } + std::cerr << "CUDA CUBLAS error: " << error << line << ":" << file << std::endl; + exit(EXIT_FAILURE); +} + +#define checkCublasErrors(error) checkCublasErrorsCall(error, __LINE__, __FILE__); + +void call_gemm(size_t dim, const half *v1_dev, const half *v2_dev, half *output_dev, cublasHandle_t blas_handle, + half &alpha, half &beta) { + checkCublasErrors(cublasHgemm(blas_handle, CUBLAS_OP_N, CUBLAS_OP_N, dim, dim, dim, + &alpha, v1_dev, dim, v2_dev, dim, &beta, output_dev, dim)); +} + +void call_gemm(size_t dim, const float *v1_dev, const float *v2_dev, float *output_dev, cublasHandle_t blas_handle, + float &alpha, float &beta) { + checkCublasErrors(cublasSgemm(blas_handle, CUBLAS_OP_N, CUBLAS_OP_N, dim, dim, dim, + &alpha, v1_dev, dim, v2_dev, dim, &beta, output_dev, dim)); +} + + +template +void run_gemm(size_t dim, real_t val_v1, real_t val_v2) { + auto size = dim * dim; + std::cout << "Testing CUBLAS with array size:" << size << std::endl; +#if USE_TENSOR_CORES == 1 + std::cout << "Using tensor cores\n"; +#else + std::cout << "Using default math\n"; +#endif + real_t *v1_dev, *v2_dev, *output_dev; + std::vector v1(size, val_v1), v2(size, val_v2), output(size); + auto byte_count = sizeof(real_t) * size; + cudaMalloc(&v1_dev, byte_count); + cudaMalloc(&v2_dev, byte_count); + cudaMalloc(&output_dev, byte_count); + cudaMemcpy(v1_dev, v1.data(), byte_count, cudaMemcpyHostToDevice); + cudaMemcpy(v2_dev, v2.data(), byte_count, cudaMemcpyHostToDevice); + cudaMemset(output_dev, 0x0, byte_count); + + cublasHandle_t blas_handle; + checkCublasErrors(cublasCreate(&blas_handle)); + checkCublasErrors(cublasSetMathMode(blas_handle, BLAS_MODE)); + + real_t alpha = 1.0f, beta = 0.0f; + call_gemm(dim, v1_dev, v2_dev, output_dev, blas_handle, alpha, beta); + + cudaMemcpy(output.data(), output_dev, byte_count, cudaMemcpyDeviceToHost); + auto expected_val = (double) dim; + for (auto i = 0; i < output.size(); i++) { + auto out = double(output[i]); + if (out != expected_val) { + std::cout << "M[" << i << "]:" << out << std::endl; + } + } + + cudaFree(v1_dev); + cudaFree(v2_dev); + cudaFree(output_dev); + std::cout << "Finished computation\n"; +} + +int main(int argc, char *argv[]) { + setbuf(stdout, nullptr); // Disable stdout buffering + auto size = 256; + run_gemm(size, 1.0f, 1.0f); + return 0; +}