Allan MacKinnon allanmac
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // | |
| // | |
| // | |
| #define WARP_SIZE 32 | |
| #define KERNEL_QUALIFIERS extern "C" __global__ | |
| #define DEVICE_INTRINSIC_QUALIFIERS __device__ __forceinline__ | |
| #define RESTRICT __restrict__ | |
| // |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // -*- compile-command: "nvcc -m 32 -arch sm_52 -Xptxas=-v,-abi=no -cubin f16.cu" ; -*- | |
| #include <stdint.h> | |
| // | |
| // | |
| // | |
| #define KERNEL_QUALIFIERS __global__ | |
| #define KERNEL_QUALIFIERS_EXTERN extern KERNEL_QUALIFIERS |
allanmac
/ int_mul.cu
Last active
November 1, 2018 04:20
Inspecting Kepler vs. Maxwell integer multiplies. Some performance insights here: https://devtalk.nvidia.com/default/topic/804281/cuda-programming-and-performance/maxwell-integer-mul-mad-instruction-counts/post/4423835/#4423835
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // -*- compile-command: "nvcc -m 32 -arch sm_50 -Xptxas=-v,-abi=no -cubin int_mul.cu" ; -*- | |
| #include <stdint.h> | |
| #include <cuda_fp16.h> | |
| // | |
| // | |
| // | |
| #define KERNEL_QUALIFIERS __global__ |
allanmac
/ scan_64.cu
Last active
August 29, 2015 14:13
Signed 64-bit reduce-add using shuffles (untested)
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // -*- compile-command: "nvcc -m 32 -arch sm_50 -Xptxas=-v,-abi=no -cubin scan_64.cu" ; -*- | |
| #include <stdint.h> | |
| // | |
| // | |
| // | |
| #define KERNEL_QUALIFIERS __global__ | |
| #define KERNEL_QUALIFIERS_EXTERN extern KERNEL_QUALIFIERS |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // -*- compile-command: "nvcc -m 32 -arch sm_50 -Xptxas=-v,-abi=no ipc.cu -o ipc" ; -*- | |
| // | |
| // | |
| // | |
| #define KERNEL_QUALIFIERS __global__ | |
| #define KERNEL_QUALIFIERS_EXTERN extern KERNEL_QUALIFIERS | |
| #define KERNEL_QUALIFIERS_EXTERN_C extern "C" KERNEL_QUALIFIERS |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // -*- compile-command: "nvcc -m 32 -arch compute_20 -Xptxas=-v,-abi=no -cubin kth.cu"; -*- | |
| #include <stdio.h> | |
| #include <stdint.h> | |
| // | |
| // | |
| // | |
| typedef uint8_t u8; |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // -*- compile-command: "nvcc -m 32 -arch sm_50 -Xptxas=-v,-abi=no -cubin xmad.cu"; -*- | |
| // | |
| // | |
| // | |
| #define KERNEL_QUALIFIERS extern "C" __global__ | |
| #define RESTRICT __restrict__ | |
| // |
allanmac
/ int_ipc.cu
Last active
August 29, 2015 14:06
Try to push an integer math kernel's IPC metric as high as possible!
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // -*- compile-command: "nvcc -m 32 -arch sm_50 -Xptxas=-v,-abi=no int_ipc.cu -o int_ipc" ; -*- | |
| // | |
| // | |
| // | |
| #define KERNEL_QUALIFIERS __global__ | |
| #define KERNEL_QUALIFIERS_EXTERN extern KERNEL_QUALIFIERS | |
| #define KERNEL_QUALIFIERS_EXTERN_C extern "C" KERNEL_QUALIFIERS |
allanmac
/ tk1_gpu_max_clock.sh
Last active
August 14, 2017 13:25
Some login tweaks for L4T 19.3 to disable USB autosuspend and lock the GPU to max MHz. I have lines 3-11 of tk1_tweaks.sh appended to my .profile.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| #!/bin/bash | |
| # | |
| # Lock GPU and MEM clocks to max MHz | |
| # | |
| # - benchmarking shows there is no need to override the MEM clock | |
| # | |
| echo | |
| echo Locking GPU clock to max MHz for debugging porpoises... | |
| # set to max |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| // -*- compile-command: "nvcc -m 32 -arch sm_50 -Xptxas=-v,-abi=no -cubin lop3.cu"; -*- | |
| #define KERNEL_QUALIFIERS extern "C" __global__ | |
| // | |
| // Bit hack: "Conditionally set or clear bits without branching" | |
| // http://graphics.stanford.edu/~seander/bithacks.html#ConditionalSetOrClearBitsWithoutBranching | |
| // | |
| // This bit hack *should* map to a single LOP3.LUT opcode: | |
| // |