froggeric/imatrix
Input files for generating the Importance Matrix Which file to use for generating the importance matrix Not all importance matrices are equal. The best results are obtained when using a source file similar to the training data. Size also matters: the bigger the model (eg: 70b vs 13b) and the higher the quant (eg: q6k_ vs iq3_xs), the bigger the source file needs to be to make an impact. Multiple input files can be combined if needed; for example: cat… See the full description on the dataset page: https://huggingface.co/datasets/froggeric/imatrix.
17293
1 move.w #$0100,d02 clr.l d13 move.w #$0400,d44 clr.l d25 move.w #$1000,d36 7NotReached:8 9 addi.b #$10,d210 add.w d0,d111 cmp.w d1,d412 bgt.s NotReached13 sub.w d2,d114 subi.w #$1000,d315 bpl.s NotReached16 move.w d1,d017 swap d018 move.w d3,d019 20 21 sep #%00100000 ;8 bit accumulator22 ldx #4 ;modifying 5 locations23 ;24 loop lda $ab1234,x ;load25 inc a ;increment26 sta $ab1234,x ;store27 dex28 bpl loop ;next29 30 31 phb ;save current data bank32 sep #%00110000 ;8 bit registers33 lda #$ab ;target bank34 pha ;push it to the stack & pull it...35 plb ;into DB, making it the default bank36 ldx #4 ;modifying 5 locations37 ;38 loop inc $1234,x ;effectively INC $AB1234,X39 dex40 bpl loop ;next41 ;42 plb ;restore previous bank43 44 sei ;IRQs off45 wai ;wait for interrupt46 lda via001 ;start of interrupt handler47 48 49 #!/bin/bash50 TARGET_DIR=("$@")51 [ "x$1" == "x" ] && TARGET_DIR=(".")52 function confirmDeletion() {53 local confirm=""54 until [ "x$confirm" == 'xy' ] || [ "x$confirm" == 'xn' ]55 do56 read -ep " Delete [y/n]: " confirm57 confirm=$(echo "$confirm" | tr [:upper:] [:lower:])58 done59 [ "x$confirm" == 'xy' ]60 }61 function deleteWithConfirmation() {62 for file in "${@}"63 do64 if rm "$file"; then65 echo " OK: $file"66 else67 echo " FAIL: $file"68 fi69 done70 }71 for i in {'*~','a.out','*.o','*.gch','*nppdf32Log*'}72 do73 echo "Files matching: $i"74 FILES=()75 while read -rd '' file76 do77 FILES+=("$file")78 echo " $file"79 done < <(find "${TARGET_DIR[@]}" -depth -iname "$i" -print0)80 if [ "x${FILES[*]}" != "x" ]; then81 if confirmDeletion; then82 deleteWithConfirmation "${FILES[@]}"83 else84 echo " Skipping"85 fi86 fi87 done88 89 90Return the open file descriptor to the calling function via the eight bit accumulator by overwriting the appropriate register stack frame element:91 92 sep #%00100000 ;select 8 bit accumulator93 lda #0 ;clear...94 xba ;.B95 lda filedes ;get file descriptor, ...96 rep #%00100000 ;select 16 bit accumulator &...97 sta reg_a,s ;overwrite .C's stack copy98 99When the accumulator is pulled it will contain the value that was in filedes.100 101Flag an error by setting the carry bit in SR:102 103 sep #%00100000 ;select 8 bit accumulator104 lda reg_sr,s ;stack copy of SR105 ora #%00000001 ;set carry bit &...106 sta reg_sr,s ;rewrite107 108Flag a successful operation by clearing the carry bit in SR:109 110 sep #%00100000 ;select 8 bit accumulator111 lda reg_sr,s ;stack copy of SR112 and #%11111110 ;clear carry bit &...113 sta reg_sr,s ;rewrite114 115class PromptFormat:116 117 botname = "Chatbort"118 username = "User"119 120 def __init__(self):121 pass122 123 #124 125 def default_system_prompt(self):126 raise NotImplementedError127 128 def first_prompt(self):129 raise NotImplementedError130 131 def subs_prompt(self):132 raise NotImplementedError133 134 def stop_conditions(self, tokenizer):135 raise NotImplementedError136 137 def encoding_options(self): # (add_bos, add_eos, encode_special_tokens)138 raise NotImplementedError139 140 def print_bot_name(self):141 return False142 143 def print_extra_newline(self):144 return False145 146 147class PromptFormat_raw(PromptFormat):148 149 description = "Model-agnostic mode simulating a raw chatlog"150 151 def __init__(self):152 super().__init__()153 pass154 155 def default_system_prompt(self):156 return \157 f"""This is a conversation between a helpful AI assistant named {self.botname} and a """ + \158 (f"""user named {self.username}.""" if self.username != "User" else """user.""")159 160 def first_prompt(self):161 return \162 f"""<|system_prompt|>\n{self.username}: <|user_prompt|>\n{self.botname}:"""163 164 def subs_prompt(self):165 return \166 f"""{self.username}: <|user_prompt|>\n{self.botname}:"""167 168 def stop_conditions(self, tokenizer):169 return \170 [self.username + ":",171 self.username[0:1] + ":",172 self.username.upper() + ":",173 self.username.lower() + ":",174 tokenizer.eos_token_id]175 176 def encoding_options(self):177 return False, False, False178 179 def print_bot_name(self):180 return True181 182########################################################183 184class PromptFormat_llama(PromptFormat):185 186 description = "Llama-chat, Llama2-chat and Mistral-instruct models"187 188 def __init__(self):189 super().__init__()190 pass191 192 def default_system_prompt(self):193 return \194 """You are a helpful, respectful and honest assistant. Always answer as helpfully as possible, while being safe. """ + \195 """Your answers should not include any harmful, unethical, racist, sexist, toxic, dangerous, or illegal content. """ + \196 """Please ensure that your responses are socially unbiased and positive in nature."""197 198 def first_prompt(self):199 return \200 """[INST] <<SYS>>\n<|system_prompt|>\n<</SYS>>\n\n<|user_prompt|> [/INST]"""201 202 def subs_prompt(self):203 return \204 """[INST] <|user_prompt|> [/INST]"""205 206 def stop_conditions(self, tokenizer):207 return \208 [tokenizer.eos_token_id]209 210 def encoding_options(self):211 return True, False, False212 213 def print_extra_newline(self):214 return True215 216 217 def build_attn_mask(self, batch_size, seq_len, past_len, input_mask, device):218 219 if input_mask is None and seq_len == 1: return None220 221 if isinstance(past_len, tuple):222 223 attn_masks = []224 225 for i in range(len(past_len[1])):226 227 attn_mask = torch.zeros((1, 1, seq_len, past_len[1][i] + seq_len), dtype = torch.float16, device = device)228 attn_mask_triu = torch.triu(torch.full((seq_len - 1, seq_len - 1), -65504.))229 attn_mask[:, :, : seq_len - 1, past_len[1][i] + 1: past_len[1][i] + seq_len] = attn_mask_triu230 231 if input_mask is not None:232 min_mask_width = min(input_mask[i].shape[-1], seq_len + past_len[1][i])233 input_mask_part = safe_move_tensor(input_mask[i][:, :min_mask_width], attn_mask.device)234 input_mask_part = input_mask_part.unsqueeze(1).unsqueeze(2)235 attn_mask[:, :, :, :min_mask_width] = torch.minimum(attn_mask[:, :, :, :min_mask_width], input_mask_part)236 237 attn_masks.append(attn_mask)238 239 return attn_masks240 241 else:242 243 attn_mask = torch.zeros((batch_size, 1, seq_len, past_len + seq_len), dtype = torch.float16, device = device)244 attn_mask_triu = torch.triu(torch.full((seq_len - 1, seq_len - 1), -65504.))245 attn_mask[:, :, : seq_len - 1, past_len + 1: past_len + seq_len] = attn_mask_triu246 247 if input_mask is not None:248 min_mask_width = min(input_mask.shape[-1], seq_len + past_len)249 input_mask_part = safe_move_tensor(input_mask[:, :min_mask_width], attn_mask.device)250 input_mask_part = input_mask_part.unsqueeze(1).unsqueeze(2)251 attn_mask[:, :, :, :min_mask_width] = torch.minimum(attn_mask[:, :, :, :min_mask_width], input_mask_part)252 253 return attn_mask254 255 256 We have to develop a Java program to calculate the sum of n natural numbers. Sum of natural number N as given as sum = 1+2+3+….+N257 258Examples:-2591+2+3+4+5 = 152601+2+3+4+5+6+7+8+9+10 = 55261 262Procedure to develop a method to find the sum of N natural numbers in Java,263 264 Take the N value.265 Declare an iterator variable and initialize it with 1 because natural numbers start with 1.266 Add iterator variable value into the sum variable267 Increase the value of the iterator variable by 1268 Repeat 3 and 4 steps until the number remains greater than the iterator variable269 270The time complexity of this procedure is O(n).271 272import java.util.Scanner;273 274public class NaturalNumberSum {275 276 // method to find sum of N natural numbers277 public static int naturalNumberSum(int number){278 279 int i = 1; // iterator variable280 // variable to store sum value281 int sum = 0;282 283 // loop to repeat the process284 while (i<=number) {285 286 // add into sum value287 sum = sum + i;288 // increase iterator variable289 i++;290 }291 292 // return sum value293 return sum;294 }295 296 public static void main(String[] args) {297 298 // declare variables299 int number = 0;300 int sum = 0;301 302 // create Scanner class object303 Scanner scan = new Scanner(System.in);304 305 // read input306 System.out.print("Enter N value:: ");307 number = scan.nextInt();308 309 // Calculate the sum value310 sum = naturalNumberSum(number);311 312 // display result313 System.out.println("Sum = "+sum);314 315 // close Scanner class objects316 scan.close();317 }318}319 320The output for different test-cases:-321 322Enter N value:: 5323Sum = 15324 325Enter N value:: 10326Sum = 55327 328In this program, we have used a while loop to find the sum of natural numbers in Java. While loop is a pre-test loop where the expression is evaluated then only statements are executed. It uses a test expression to control the loop. Before every iteration of the loop, the test expression is evaluated.329 330Also See:-331 332 Sum of digits of a number333 The sum of even digits in a number334 Sum of odd digits in a number335 Sum of first & last digit of a number336 The Sum of Digits Until Single Digit337 338We can also use for loop instead of using a while loop. The for loop is also a pre-test loop, where first of all initialization expression is evaluated then the condition is checked and if the condition is true then only the statements of the for loop are executed.339 340public static int naturalNumberSum(int number){341 342 int sum = 0;343 344 for(int i=1; i<=number; i++)345 sum+=i;346 347 return sum;348 }349 350Or,351 352public static int naturalNumberSum(int number){353 354 int sum = 0;355 356 for(int i=1; ; sum+=i, i++)357 if(i>number) return sum;358 359}360 361The time complexity of all above methods are O(n).362Sum of Natural Numbers in Java without using the loop363 364We can also do the same work without using the loop. The formula for this operation,365 366Sum = n * (n+1) / 2;367 368Example:-369Sum of first 10 natural numbers = 10*(10+1)/2 = 10*11/2 = 5*11 = 55370 371It is the best way to find the sum of natural numbers. The time complexity of this method is O(1).372 373import java.util.Scanner;374 375public class NaturalNumberSum {376 377 // method to find sum of N natural numbers378 public static int naturalNumberSum(int number){379 return number*(number+1)/2;380 }381 382 public static void main(String[] args) {383 384 // declare variables385 int number = 0;386 int sum = 0;387 388 // create Scanner class object389 Scanner scan = new Scanner(System.in);390 391 // read input392 System.out.print("Enter N value:: ");393 number = scan.nextInt();394 395 // Calculate the sum value396 sum = naturalNumberSum(number);397 398 // display result399 System.out.println("Sum = "+sum);400 401 // close Scanner class objects402 scan.close();403 }404}405 406Using recursion407 408We already developed java program to find the sum of the natural number using for loop, or while loop, or without using the loop. Now we will find the same using the recursion technique. In Recursion, We divided the big problems into smaller problems or sub-problems.409 410Sum of N natural numbers given as 1+2+….+(n-1)+n. So, the problem can be divided as n + ( (n-1) +… + 2 + 1 )411 412General case for finding the sum of natural number => sum(n) = n + sum(n-1); Similarly, the base case for finding the sum of natural number => sum(0) = 0; or sum(1) = 1;413 414import415 416module jtransmissiongatetb;417 wire y;418 reg a,control;419 jtransmissiongate jtgate(y,control,a);420 initial421 begin422 $display ("RESULT\ta\ty");423 424 a = 0; control = 0; # 50; // Initial value is set425 if ( y === 1'bz ) // Test for inversion426 $display ("PASS \t%d\t%d",a,y);427 else428 $display ("FAIL \t%d\t%d",a,y);429 control = 1; # 50; // Simply change the control signal430 control = 0; # 50; // Simply change the control signal431 control = 1; # 50; // Simply change the control signal432 control = 0; # 50; // Simply change the control signal433 434 a = 0; control = 1; # 50; // Initial value is set435 if ( y === 0 ) // Test for inversion436 $display ("PASS \t%d\t%d",a,y);437 else438 $display ("FAIL \t%d\t%d",a,y);439 440 a = 1; control = 0; # 50; // Another value441 if ( y === 1'bz ) // Test for inversion442 $display ("PASS \t%d\t%d",a,y);443 else444 $display ("FAIL \t%d\t%d",a,y);445 control = 1; # 50; // Simply change the control signal446 control = 0; # 50; // Simply change the control signal447 control = 1; # 50; // Simply change the control signal448 control = 0; # 50; // Simply change the control signal449 450 a = 1; control = 1; # 50; // Another value451 if ( y === 1 ) // Test for inversion452 $display ("PASS \t%d\t%d",a,y);453 else454 $display ("FAIL \t%d\t%d",a,y);455 456 end457 //enabling the wave dump458 initial begin459 $dumpfile("dump.vcd"); $dumpvars;460 end461endmodule462 463 464module jtransmissiongate(y,control,a);465 output y;466 input a,control;467 468 wire cbar;469 470 assign cbar = ~control;471 472 nmos n1(y,a,control);473 pmos p1(y,a,cbar);474 //cmos c1(y,a,control,cbar);475 476endmodule477 478 479module juniversalShiftRegisterTb;480 wire [3:0] DATAOUT;481 reg clock, reset;482 reg [1:0] MODE;483 reg [3:0] DATAIN;484 485 juniversalShiftRegister jusr(DATAOUT, clock, reset, MODE, DATAIN);486 487 initial488 begin489 clock =0; MODE = 2'b00; DATAIN = 4'b0000;490 reset = 1; #10; reset = 0; #10;491 492 $display("RSLT\tD == DOUT");493 // Start testing Right Shift mode494 MODE = 2'b00; reset = 1; #10; reset = 0; #10;495 MODE = 2'b01; DATAIN = 4'b0011; #10;496 if ( DATAOUT === 4'b1000 ) // look at previous value of DATAOUT as well497 $display("PASS\t%p is %p with %p", DATAIN, MODE, DATAOUT);498 else499 $display("FAIL\t%p is %p with %p", DATAIN, MODE, DATAOUT);500 MODE = 2'b01; DATAIN = 4'b0011; #10;501 if ( DATAOUT === 4'b1100 ) // look at previous value of DATAOUT as well502 $display("PASS\t%p is %p with %p", DATAIN, MODE, DATAOUT);503 else504 $display("FAIL\t%p is %p with %p", DATAIN, MODE, DATAOUT);505 506 507 // Start testing Left Shift mode508 MODE = 2'b00; reset = 1; #10; reset = 0; #10;509 MODE = 2'b10; DATAIN = 4'b0111; #10;510 if ( DATAOUT === 4'b0001 ) //511 $display("PASS\t%p is %p with %p", DATAIN, MODE, DATAOUT);512 else513 $display("FAIL\t%p is %p with %p", DATAIN, MODE, DATAOUT);514 MODE = 2'b10; DATAIN = 4'b0111; #10;515 if ( DATAOUT === 4'b0011 ) //516 $display("PASS\t%p is %p with %p", DATAIN, MODE, DATAOUT);517 else518 $display("FAIL\t%p is %p with %p", DATAIN, MODE, DATAOUT);519 520 521 // Start testing parallel load mode522 MODE = 2'b00; reset = 1; #10; reset = 0; #10;523 MODE = 2'b11; DATAIN = 4'b1010; #10;524 if ( DATAOUT === 4'b1010 )525 $display("PASS\t%p is %p with %p", DATAIN, MODE, DATAOUT);526 else527 $display("FAIL\t%p is %p with %p", DATAIN, MODE, DATAOUT);528 529 #20;530 $finish;531 end532 533 534 //enabling the wave dump535 initial begin536 $dumpfile("dump.vcd"); $dumpvars;537 end538 539 540 always #5 clock = ~clock;541 542endmodule543 544#!/bin/bash545# use predefined variables to access passed arguments546#echo arguments to the shell547echo $1 $2 $3 ' -> echo $1 $2 $3'548 549# We can also store arguments from bash command line in special array550args=("$@")551#echo arguments to the shell552echo ${args[0]} ${args[1]} ${args[2]} ' -> args=("$@"); echo ${args[0]} ${args[1]} ${args[2]}'553 554#use $@ to print out all arguments at once555echo $@ ' -> echo $@'556 557# use $# variable to print out558# number of arguments passed to the bash script559echo Number of arguments passed: $# ' -> echo Number of arguments passed: $#'560 561Let’s try executing this script and providing three arguments.562 563$ ./arguments.sh Bash Scripting Tutorial564 565The results when we execute this script:566 567Bash Scripting Tutorial -> echo $1 $2 $3568Bash Scripting Tutorial -> args=("$@"); echo ${args[0]} ${args[1]} ${args[2]}569Bash Scripting Tutorial -> echo $@570Number of arguments passed: 3 -> echo Number of arguments passed: $#571 572Executing shell commands with bash573 574The best way to execute a separate shell command inside of a Bash script is by creating a new subshell through the $( ) syntax. Check the example below where we echo the result of running the uname -o command.575 576#!/bin/bash577# use a subshell $() to execute shell command578echo $(uname -o)579# executing bash command without subshell580echo uname -o581 582Notice that in the final line of our script, we do not execute the uname command within a subshell, therefore the text is taken literally and output as such.583 584$ uname -o585GNU/LINUX586$ ./subshell.sh587GNU/LINUX588uname -o589 590Reading User Input591 592We can use the read command to read input from the user. This allows a user to interact with a Bash script and help dictate the way it proceeds. Here’s an example:593 594#!/bin/bash595 596echo -e "Hi, please type the word: \c "597read word598echo "The word you entered is: $word"599echo -e "Can you please enter two words? "600read word1 word2601echo "Here is your input: \"$word1\" \"$word2\""602echo -e "How do you feel about bash scripting? "603# read command now stores a reply into the default build-in variable $REPLY604read605echo "You said $REPLY, I'm glad to hear that! "606echo -e "What are your favorite colours ? "607# -a makes read command to read into an array608read -a colours609echo "My favorite colours are also ${colours[0]}, ${colours[1]} and ${colours[2]}:-)"610 611Our Bash script asks multiple questions and then is able to repeat the information back to us through variables and arrays:612 613$ ./read.sh614Hi, please type the word: Linuxconfig.org615The word you entered is: Linuxconfig.org616Can you please enter two words?617Debian Linux618Here is your input: "Debian" "Linux"619How do you feel about bash scripting?620good621You said good, I'm glad to hear that!622What are your favorite colours ?623blue green black624My favorite colours are also blue, green and black:-)625 626Bash Trap Command627 628The trap command can be used in Bash scripts to catch signals sent to the script and then execute a subroutine when they occur. The script below will detect a Ctrl + C interrupt.629 630#!/bin/bash631# bash trap command632trap bashtrap INT633# bash clear screen command634clear;635# bash trap function is executed when CTRL-C is pressed:636# bash prints message => Executing bash trap subrutine !637bashtrap()638{639 echo "CTRL+C Detected !...executing bash trap !"640}641# for loop from 1/10 to 10/10642for a in `seq 1 10`; do643 echo "$a/10 to Exit."644 sleep 1;645done646echo "Exit Bash Trap Example!!!"647 648In the output below you can see that we try to Ctrl + C two times but the script continues to execute.649 650$ ./trap.sh6511/10 to Exit.6522/10 to Exit.653^CCTRL+C Detected !...executing bash trap !6543/10 to Exit.6554/10 to Exit.6565/10 to Exit.6576/10 to Exit.6587/10 to Exit.659^CCTRL+C Detected !...executing bash trap !6608/10 to Exit.6619/10 to Exit.66210/10 to Exit.663Exit Bash Trap Example!!!664 665Arrays666 667Bash is capable of storing values in arrays. Check the sections below for two different examples.668Declare simple bash array669 670This example declares an array with four elements.671 672#!/bin/bash673#Declare array with 4 elements674ARRAY=( 'Debian Linux' 'Redhat Linux' Ubuntu Linux )675# get number of elements in the array676ELEMENTS=${#ARRAY[@]}677 678# echo each element in array679# for loop680for (( i=0;i<$ELEMENTS;i++)); do681 echo ${ARRAY[${i}]}682done683 684Executing the script will output the elements of our array:685 686$ ./arrays.sh687Debian Linux688Redhat Linux689Ubuntu690Linux691 692Read file into bash array693 694Rather than filling out all of the elements of our array in the Bash script itself, we can program our script to read input and put it into an array.695 696#!/bin/bash697# Declare array698declare -a ARRAY699# Link filedescriptor 10 with stdin700exec 10<&0701# stdin replaced with a file supplied as a first argument702exec < $1703let count=0704 705while read LINE; do706 707 ARRAY[$count]=$LINE708 ((count++))709done710 711echo Number of elements: ${#ARRAY[@]}712# echo array's content713echo ${ARRAY[@]}714# restore stdin from filedescriptor 10715# and close filedescriptor 10716exec 0<&10 10<&-717 718Now let’s execute the script and store four elements in the array by using a file’s contents for input.719 720$ cat bash.txt721Bash722Scripting723Tutorial724Guide725$ ./bash-script.sh bash.txt726Number of elements: 4727Bash Scripting Tutorial Guide728 729 730 731 732 *** CompressADPCM3 ***733 734 ; JoinCode = CompressADPCM3(Source, Length, Destination, JoinCode)735 ; d0 a0 d0 a1 d1736 ;737 ; This function compresses a RAW sample to a given memory. The738 ; result is a 3bit ADPCM code. The destination buffer must be739 ; at least (Length+7)/8*3 bytes in size.740 ;741 ; Function of the JoinCode: See above.742 743 XDEF _CompressADPCM3744_CompressADPCM3745 movem.l d2-d4,-(sp)746 747 move.w d1,d3 ; d3=EstMax748 swap d1749 move.w d1,d2 ; d2=Delta750 bne.s c3_loop751 moveq #5,d2752 753c3_loop moveq #0,d1 ; d1=Shifter754 bsr.s c3_byte755 lsl.b #3,d1756 bsr.s c3_byte757 lsl.w #3,d1758 bsr.s c3_byte759 lsl.w #3,d1760 bsr.s c3_byte761 lsl.w #3,d1762 bsr.s c3_byte763 lsl.l #3,d1764 bsr.s c3_byte765 lsl.l #3,d1766 bsr.s c3_byte767 lsl.l #3,d1768 bsr.s c3_byte769 swap d1770 move.b d1,(a1)+771 rol.l #8,d1772 move.b d1,(a1)+773 rol.l #8,d1774 move.b d1,(a1)+775 776 subq.l #8,d0 ; d0=Counter777 bhi.s c3_loop778 779 move.w d2,d0 ; -> d0=JoinCode780 swap d0781 move.w d3,d0782 783 movem.l (sp)+,d2-d4784 rts785 786c3_byte move.b (a0)+,d4787 ext.w d4788 asl.w #6,d4789 sub.w d3,d4790 bpl.s c3_positive791 or.b #%100,d1792 neg.w d4793c3_positive sub.w d2,d4794 bls.s c3_00795 sub.w d2,d4796 bls.s c3_01797 sub.w d2,d4798 bls.s c3_10799c3_11 or.b #%11,d1800 bra.s c3_00801c3_10 or.b #%10,d1802 bra.s c3_00803c3_01 or.b #%01,d1804c3_00 bsr.s adaptive805 rts806 807 808 809 *** Adaptions-Routine ***810 811adaptive ; d1 = SignBit + DataBit812 813 move.w d2,d4814 lsr.w #1,d4815 btst #1,d1816 beq.s d3_0817d3_1 btst #0,d1818 beq.s d3_10819d3_11 add.w d2,d4820 add.w d2,d4821 add.w d2,d4822 mulu #$6607,d2823 bra.s d3_sign824d3_10 add.w d2,d4825 add.w d2,d4826 mulu #$4D14,d2827 bra.s d3_sign828d3_0 btst #0,d1829 beq.s d3_00830d3_01 add.w d2,d4831 mulu #$3A9F,d2832 bra.s d3_sign833d3_00 mulu #$399A,d2834d3_sign btst #2,d1835 beq.s d3_add836 neg.w d4837d3_add add.w d4,d3838 add.l #8192,d2839 moveq #14,d4840 asr.l d4,d2841 rts842 843 844 END845 846 mov847 848 849 mov bp, 255850 851traceloop:852 853; Evaluate whether sample point is inside or outside the shape:854;855; ( a & ( b | c ) ) | ( b & c ) = 0 <=> voxel overlaps fractal856 857 push bx858 859 mov dx, bx860 or dx, cx861 and dx, ax862 and bx, cx863 or dx, bx864 865 pop bx866 867; Ignore the lower bits or the fractal will be too fine to see868 869 shr dx, 6870 jz endtrace871 872 dec bp873 jnz traceloop874 875endtrace:876 877; BP is 255 - the distance we had to trace878 879 mov dx, bp880 not dl881 882; Plot pixel883 884 mov ds:[di],dl885 inc di886 887 888 889 890 891 892// ******************893 894#include <torch/extension.h>895#include <c10/cuda/CUDAGuard.h>896#include <ATen/cuda/CUDAContext.h>897#include <cuda_runtime.h>module jtransmissiongatetb;898 wire y;899 reg a,control;900 jtransmissiongate jtgate(y,control,a);901 initial902 begin903 $display ("RESULT\ta\ty");904 905 a = 0; control = 0; # 50; // Initial value is set906 if ( y === 1'bz ) // Test for inversion907 $display ("PASS \t%d\t%d",a,y);908 else909 $display ("FAIL \t%d\t%d",a,y);910 control = 1; # 50; // Simply change the control signal911 control = 0; # 50; // Simply change the control signal912 control = 1; # 50; // Simply change the control signal913 control = 0; # 50; // Simply change the control signal914 915 a = 0; control = 1; # 50; // Initial value is set916 if ( y === 0 ) // Test for inversion917 $display ("PASS \t%d\t%d",a,y);918 else919 $display ("FAIL \t%d\t%d",a,y);920 921 a = 1; control = 0; # 50; // Another value922 if ( y === 1'bz ) // Test for inversion923 $display ("PASS \t%d\t%d",a,y);924 else925 $display ("FAIL \t%d\t%d",a,y);926 control = 1; # 50; // Simply change the control signal927 control = 0; # 50; // Simply change the control signal928 control = 1; # 50; // Simply change the control signal929 control = 0; # 50; // Simply change the control signal930 931 a = 1; control = 1; # 50; // Another value932 if ( y === 1 ) // Test for inversion933 $display ("PASS \t%d\t%d",a,y);934 else935 $display ("FAIL \t%d\t%d",a,y);936 937 end938 //enabling the wave dump939 initial begin940 $dumpfile("dump.vcd"); $dumpvars;941 end942endmodule943#include <cuda_fp16.h>944#include <cstdint>945#include <cstdio>946#include <pybind11/pybind11.h>947#include <pybind11/stl.h>948 949#include "config.h"950 951#include "cuda/pack_tensor.cuh"952#include "cuda/quantize.cuh"953#include "cuda/q_matrix.cuh"954#include "cuda/q_attn.cuh"955#include "cuda/q_mlp.cuh"956#include "cuda/q_gemm.cuh"957#include "cuda/rms_norm.cuh"958#include "cuda/rope.cuh"959#include "cuda/cache.cuh"960#include "cuda/h_gemm.cuh"961 962#include "cpp/quantize_func.h"963#include "cpp/sampling.h"964 965#include "cpp/util.h"966 967// Some decluttering macros968 969#define TORCH_CHECK_DTYPE(__x, __dtype) TORCH_CHECK((__x).dtype() == torch::__dtype, #__x " is incorrect datatype, must be " #__dtype)970#define TORCH_CHECK_DTYPE_OPT(__x, __dtype) TORCH_CHECK((__x).device().is_meta() || (__x).dtype() == torch::__dtype, #__x " is incorrect datatype, must be " #__dtype)971#define TORCH_CHECK_SHAPES(__x, __dim_x, __y, __dim_y, __scale_y) TORCH_CHECK((__x).size(__dim_x) == (__y).size(__dim_y) * __scale_y, #__x " and " #__y " have incompatible shapes")972#define TORCH_CHECK_SHAPES_OPT(__x, __dim_x, __y, __dim_y, __scale_y) TORCH_CHECK((__x).device().is_meta() || (__x).size(__dim_x) == (__y).size(__dim_y) * __scale_y, #__x " and " #__y " have incompatible shapes")973 974 975// Packing functions976 977void pack_rows_4978(979 torch::Tensor input,980 torch::Tensor output981)982{983 const at::cuda::OptionalCUDAGuard device_guard(device_of(input));984 985 TORCH_CHECK_DTYPE(input, kShort);986 TORCH_CHECK_DTYPE(output, kInt);987 TORCH_CHECK_SHAPES(input, 0, output, 0, 1);988 TORCH_CHECK_SHAPES(input, 1, output, 1, 8);989 990 int rows = input.size(0);991 int columns = input.size(1);992 993 pack_rows_4_cuda994 (995 (uint16_t*) input.data_ptr(),996 (uint32_t*) output.data_ptr(),997 rows,998 columns999 );1000}1001 1002void pack_columns1003(1004 torch::Tensor input,1005 torch::Tensor output,1006 int bits1007)1008{1009 const at::cuda::OptionalCUDAGuard device_guard(device_of(input));1010 1011 TORCH_CHECK_DTYPE(input, kShort);1012 TORCH_CHECK_DTYPE(output, kInt);1013 TORCH_CHECK_SHAPES(input, 1, output, 1, 1);1014 1015 int in_rows = input.size(0);1016 int columns = input.size(1);1017 int out_rows = output.size(0);1018 int exp_out_rows = in_rows * bits / 32;1019 TORCH_CHECK(out_rows == exp_out_rows, "Wrong output shape for input and bitrate")1020 1021 pack_columns_cuda1022 (1023 (uint16_t*) input.data_ptr(),1024 (uint32_t*) output.data_ptr(),1025 in_rows,1026 out_rows,1027 columns,1028 bits1029 );1030}1031 1032#include "quantize_func.h"1033#include "../cuda/quantize.cuh"1034 1035void quantize_range1036(1037 torch::Tensor quant,1038 torch::Tensor scale,1039 torch::Tensor out_q,1040 float qzero,1041 float maxq,1042 torch::Tensor hessian_inv,1043 torch::Tensor weights,1044 torch::Tensor error,1045 int a,1046 int b1047)1048{1049 int columns = weights.size(1);1050 int hcolumns = hessian_inv.size(1);1051 1052 for (int c = a; c < b; c++)1053 {1054 quantize_cuda1055 (1056 ((const float*) weights.data_ptr()) + c * columns,1057 ((float*) quant.data_ptr()) + c * columns,1058 (const float*) scale.data_ptr(),1059 out_q.device().is_meta() ? NULL : ((uint16_t*) out_q.data_ptr()) + c * columns,1060 1,1061 columns,1062 qzero,1063 maxq1064 );1065 1066 adjust_error_row_cuda1067 (1068 (const float*) hessian_inv.data_ptr(),1069 (float*) error.data_ptr(),1070 (const float*) weights.data_ptr(),1071 (const float*) quant.data_ptr(),1072 c,1073 columns,1074 hcolumns1075 );1076 1077 vv_mul_sub_cuda1078 (1079 ((const float*) hessian_inv.data_ptr()) + c * hcolumns + c,1080 ((const float*) error.data_ptr()) + c * columns,1081 ((float*) weights.data_ptr()) + c * columns,1082 b - c,1083 columns1084 );1085 }1086 1087 torch::Tensor x = hessian_inv.slice(0, a, b).slice(1, b).transpose(0, 1);1088 torch::Tensor y = error.slice(0, a, b);1089 weights.slice(0, b).addmm_(x, y, 1.0f, -1.0f);1090}1091 1092//---------------------------------------------------------------------------------------------------1093//---------------------------------------------------------------------------------------------------1094 1095__forceinline__ __device__ half dot22_8_h(half2(&dq)[4], const half* a_ptr, const half g_result, const half qs_h)1096{1097 half2 result = {};1098 const half2* a2_ptr = (const half2*)a_ptr;1099 #pragma unroll1100 for (int i = 0; i < 4; i++) result = __hfma2(dq[i], *a2_ptr++, result);1101 half result_h = __hadd(__low2half(result), __high2half(result));1102 return __hfma(result_h, qs_h, g_result);1103}1104 1105__forceinline__ __device__ half dot22_16_h(half2(&dq)[8], const half* a_ptr, const half g_result, const half qs_h)1106{1107 half2 result = {};1108 const half2* a2_ptr = (const half2*)a_ptr;1109 #pragma unroll1110 for (int i = 0; i < 8; i++) result = __hfma2(dq[i], *a2_ptr++, result);1111 half result_h = __hadd(__low2half(result), __high2half(result));1112 return __hfma(result_h, qs_h, g_result);1113}1114 1115__forceinline__ __device__ half dot22_32_h(half2(&dq)[16], const half* a_ptr, const half g_result, const half qs_h)1116{1117 half2 result = {};1118 const half2* a2_ptr = (const half2*)a_ptr;1119 #pragma unroll1120 for (int i = 0; i < 16; i += 1) result = __hfma2(dq[i], *a2_ptr++, result);1121 half result_h = __hadd(__low2half(result), __high2half(result));1122 return __hfma(result_h, qs_h, g_result);1123}1124 1125name: Build Wheels1126 1127on: workflow_dispatch1128 1129jobs:1130 build_wheels:1131 name: ${{ matrix.os }} Python ${{ matrix.pyver }} CUDA ${{ matrix.cuda }}1132 runs-on: ${{ matrix.os }}1133 strategy:1134 matrix:1135 os: [ubuntu-20.04, windows-latest]1136 pyver: ["3.8", "3.9", "3.10", "3.11"]1137 cuda: ["11.7.0", "11.8.0", "12.1.1"]1138 defaults:1139 run:1140 shell: pwsh1141 env:1142 CUDAVER: ${{ matrix.cuda }}1143 PYVER: ${{ matrix.pyver }}1144 1145 steps:1146 - name: Free Disk Space1147 uses: jlumbroso/free-disk-space@v1.2.01148 if: runner.os == 'Linux'1149 with:1150 tool-cache: false1151 android: true1152 dotnet: true1153 haskell: true1154 large-packages: false1155 swap-storage: false1156 1157 - uses: actions/checkout@v31158 - uses: actions/setup-python@v31159 with:1160 python-version: ${{ matrix.pyver }}1161 1162 - name: Setup Mamba1163 uses: conda-incubator/setup-miniconda@v2.2.01164 with:1165 activate-environment: "build"1166 python-version: ${{ matrix.pyver }}1167 miniforge-variant: Mambaforge1168 miniforge-version: latest1169 use-mamba: true1170 add-pip-as-python-dependency: true1171 auto-activate-base: false1172 1173 - name: Install Dependencies1174 run: |1175 $cudaVersion = $env:CUDAVER1176 $cudaVersionPytorch = $env:CUDAVER.Remove($env:CUDAVER.LastIndexOf('.')).Replace('.','')1177 1178 $cudaChannels = ''1179 $cudaNum = [int]$cudaVersion.substring($cudaVersion.LastIndexOf('.')+1)1180 while ($cudaNum -ge 0) { $cudaChannels += '-c nvidia/label/cuda-' + $cudaVersion.Remove($cudaVersion.LastIndexOf('.')+1) + $cudaNum + ' '; $cudaNum-- }1181 mamba install -y 'cuda' $cudaChannels.TrimEnd().Split()1182 1183 if (!(mamba list cuda)[-1].contains('cuda')) {sleep -s 10; mamba install -y 'cuda' $cudaChannels.TrimEnd().Split()}1184 if (!(mamba list cuda)[-1].contains('cuda')) {throw 'CUDA Toolkit failed to install!'}1185 1186 if ([version]$env:CUDAVER -lt [version]'11.8.0') {$torchver = "torch==2.0.1"} else {$torchver = "torch==2.1.0"}1187 python -m pip install $torchver --index-url https://download.pytorch.org/whl/cu$cudaVersionPytorch1188 1189 python -m pip install build wheel safetensors sentencepiece ninja1190 1191 - name: Build Wheel1192 run: |1193 $env:CUDA_PATH = $env:CONDA_PREFIX1194 $env:CUDA_HOME = $env:CONDA_PREFIX1195 1196 $cudaVersion = $env:CUDAVER1197 $cudaVersionPytorch = $env:CUDAVER.Remove($env:CUDAVER.LastIndexOf('.')).Replace('.','')1198 $BUILDTAG = "+cu$cudaVersionPytorch"1199 1200 if ($IsLinux) {$env:LD_LIBRARY_PATH = $env:CONDA_PREFIX + '/lib:' + $env:LD_LIBRARY_PATH}