Allan MacKinnon allanmac
This file contains 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_35 -Xptxas=-v,-abi=no -cubin short4.cu"; -*- | |
| #include <stdint.h> | |
| typedef uint32_t u32; | |
| typedef uint64_t u64; | |
| typedef union | |
| { | |
| short4 s16v4; |
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 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 |
This file contains 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__ | |
| // |
This file contains 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 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 |
allanmac
/ scan_64.cu
Last active
August 29, 2015 14:13
Signed 64-bit reduce-add using shuffles (untested)
This file contains 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 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 |
This file contains 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 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
| #include <stdio.h> | |
| // | |
| // | |
| // | |
| static | |
| void | |
| cuda_assert(const cudaError_t code, const char* const file, const int line, const bool abort) | |
| { |
This file contains 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
| __global__ | |
| void fmuladdTest(float* const values) | |
| { | |
| const unsigned int tidx = threadIdx.x; | |
| const float b = values[ tidx]; | |
| float a = values[2*tidx]; | |
| a = __fmul_rn(a, b); | |
| a = __fadd_rn(a, 0.73f); |
OlderNewer