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
| __global__ | |
| void fmaTest(float* const values) | |
| { | |
| const unsigned int tidx = threadIdx.x; | |
| const float b = values[ tidx]; | |
| float a = values[2*tidx]; | |
| a = __fmaf_rn(a, b, 0.73f); | |
| a = __fmaf_rn(a, b, 0.37f); |
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 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> | |
| // | |
| // | |
| // | |
| #define DEVICE_INTRINSIC_QUALIFIERS __device__ __forceinline__ | |
| DEVICE_INTRINSIC_QUALIFIERS | |
| unsigned int |
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> | |
| // | |
| // | |
| // | |
| #define TYPE unsigned int | |
| #define REPS 1 |
allanmac
/ natural.cu
Last active
December 14, 2015 06:38
The `setp` and `selp` instructions are your friends.
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
| extern "C" | |
| __global__ | |
| void natural(const unsigned int b, | |
| const unsigned int c, | |
| const unsigned int y, | |
| const unsigned int z, | |
| const unsigned int id, | |
| unsigned int* const out) | |
| { | |
| const bool flag = (id == 1); |
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 KERNEL_QUALIFIERS extern "C" __global__ | |
| KERNEL_QUALIFIERS | |
| void fdimfTest(const float x, const float y, float* const fout) | |
| { | |
| fout[threadIdx.x] = fdimf(x,y); | |
| } | |
| KERNEL_QUALIFIERS | |
| void fdimfTest2(const float x, const float y, float* const fout) |
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 KERNEL_QUALIFIERS extern "C" __global__ | |
| KERNEL_QUALIFIERS | |
| void shflmax(const int* const vin, int* const vout) | |
| { | |
| int v = vin[threadIdx.x]; | |
| v = max(v,__shfl_xor(v,16)); | |
| v = max(v,__shfl_xor(v, 8)); |
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 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 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 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() \ |