Allan MacKinnon allanmac
allanmac
/ alpha_gamma_test.svg
Last active
February 22, 2016 02:31
See here: http://www.realtimerendering.com/blog/a-png-puzzle/
Loading
Sorry, something went wrong. Reload?
Sorry, we cannot display this file.
Sorry, this file is invalid so it cannot be displayed.
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 -D THRUST_SORT_TYPE=uint32_t -arch sm_50 -o sort sort_32.cu"; -*- | |
| #include <thrust/host_vector.h> | |
| #include <thrust/device_vector.h> | |
| #include <thrust/generate.h> | |
| #include <thrust/sort.h> | |
| #include <thrust/copy.h> | |
| #include <algorithm> | |
| #include <cstdlib> |
allanmac
/ probe_bw.cu
Last active
October 10, 2019 15:27
Measure achieved bandwidth when performing 128, 256 or 512 byte transactions on a multi-megabyte extent. This appears to reproduce @Genoil's original findings: https://devtalk.nvidia.com/default/topic/878455/cuda-programming-and-performance/gtx750ti-and-buffers-gt-1gb-on-win7
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 64 -arch compute_30 -Xptxas=-v -o probe_bw probe_bw.cu"; -*- | |
| // | |
| // Copyright 2015 Allan MacKinnon <allanmac@alum.mit.edu> | |
| // | |
| // Permission is hereby granted, free of charge, to any person obtaining | |
| // a copy of this software and associated documentation files (the | |
| // "Software"), to deal in the Software without restriction, including | |
| // without limitation the rights to use, copy, modify, merge, publish, | |
| // distribute, sublicense, and/or sell copies of the Software, and to |
allanmac
/ assert_cuda.c
Last active
April 11, 2026 23:42
A tiny example of CUDA + OpenGL interop with write-only surfaces and CUDA kernels. Uses GLFW+GLAD.
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
| // | |
| // | |
| // | |
| #include <stdlib.h> | |
| #include <stdio.h> | |
| // | |
| // | |
| // |
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; |