Allan MacKinnon allanmac
allanmac
/ bfe64.cu
Last active
December 16, 2015 23:29
What is the best way to extract up to 32 bits that straddle the 32-bit boundary of a 64-bit word given a constant starting position and number of bits? On sm_35 the SHF.R.CLAMP opcode can accomplish this in two instructions. For sm_12-sm_30 devices as many as four instructions are required.
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 <stdio.h> | |
| // | |
| // | |
| // | |
| #define DEVICE_INTRINSIC_QUALIFIERS __device__ __forceinline__ | |
| // | |
| // |
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
| all: | |
| nvcc -m 32 -Xptxas=-v,-abi=no \ | |
| -gencode=arch=compute_11,code=sm_11 \ | |
| -gencode=arch=compute_12,code=sm_12 \ | |
| -gencode=arch=compute_20,code=sm_21 \ | |
| -gencode=arch=compute_30,code=sm_30 \ | |
| -gencode=arch=compute_35,code=sm_35 \ | |
| blocks.cu -o blocks |
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 KERNEL_QUALIFIERS extern "C" __global__ | |
| // | |
| // | |
| // | |
| #define REPEAT1() \ | |
| REPS(0) | |
| #define REPEAT4() \ |
allanmac
/ sync.cu
Created
July 14, 2013 21:21
Examine the SASS that's generated for barrier reduction operations: __syncthreads_count(), __syncthreads_or(), __syncthreads_and() as well as the regular __syncthreads() barrier op. Somewhat surprisingly these are not mapped to a number of SASS ops. The barrier reductions are executed and the result is moved from a "barrier register" to a regula…
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 KERNEL_QUALIFIERS extern "C" __global__ | |
| // | |
| // | |
| // |
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 KERNEL_QUALIFIERS __global__ | |
| #define VOLATILE volatile | |
| #define DEVICE_INTRINSIC_QUALIFIERS __device__ __forceinline__ | |
| #define DEVICE_STATIC_FUNCTION_QUALIFIERS static DEVICE_FUNCTION_QUALIFIERS | |
| #define DEVICE_STATIC_INTRINSIC_QUALIFIERS static DEVICE_INTRINSIC_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
| #include <stdio.h> | |
| // | |
| // | |
| // | |
| #define WARP_SIZE 32 | |
| #define VOLATILE volatile |
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 KERNEL_QUALIFIERS extern "C" __global__ | |
| // | |
| // | |
| // |
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 <stdio.h> | |
| #include <cuda.h> | |
| int main(int argc, char** argv) | |
| { | |
| cuInit(0); | |
| int count; |
allanmac
/ warp_scan.inl
Last active
December 25, 2015 00:28
The macro at the bottom of "warp_scan.inl" is used to declare an optimal CUDA warp scan primitive without using C++ templates and specialization. The macro supports 32-bit PTX types (u32/s32/f32) and can generate inclusive and exclusive scans over any appropriate PTX two-argument operator (add,sub,min,max,mul,div,rem,etc). See examples below.
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
| #pragma once | |
| // | |
| // | |
| // | |
| #define PXL_WARP_SCAN_SHFL(_op,_vT,_opT,_regC,_exc,_exc0,_excP) | |
| //////////////////////////////////////////////////////////////////////// | |
| // |
allanmac
/ float3 SoA to AoS
Last active
June 30, 2018 16:12
A strategy for converting a float3 SoA into AoS without using shared memory.
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
| =============================================================================================== | |
| Load three arrays (x, y and z) in SoA order, repack them and store them in AoS order. | |
| Strategy: each warp permutes its load lane with: | |
| (rowNum + (laneId() * 3)) & 31 | |
| This will convert SoA into AoS but with x/y/z staggered across rows of registers. |