Device intrinsic host fallbacks

Cache-hinted loads/stores are plain pointer access in the one-address-space sim; the hint (streaming / last-use / global) has no effect on correctness.

Defined in code/simulator.h.

This page documents 29 symbols: 29 functions.

Functions

Signature	Description	Availability	Location	Example
inline float __cosf(float x)		sim only	simulator.h:792
inline float __exp10f(float x)		sim only	simulator.h:787
inline float __expf(float x)		sim only	simulator.h:786	View
inline float __fdividef(float a, float b)		sim only	simulator.h:795
inline T __ldca(const T* p)		sim only	simulator.h:764
inline T __ldcg(const T* p)		sim only	simulator.h:766
inline T __ldcs(const T* p)		sim only	simulator.h:765	View
inline T __ldg (const T* p)	Cache-hinted loads/stores are plain pointer access in the one-address-space sim; the hint (streaming / last-use / global) has no effect on correctness.	sim only	simulator.h:763	View
inline T __ldlu(const T* p)		sim only	simulator.h:767
inline float __log10f(float x)		sim only	simulator.h:790
inline float __log2f(float x)		sim only	simulator.h:789
inline float __logf(float x)		sim only	simulator.h:788
inline int __mul24 (int a, int b)		sim only	simulator.h:797	View
inline int __mulhi (int a, int b)		sim only	simulator.h:799
inline void __nanosleep(unsigned )	Memory fences: the sim runs one fiber at a time and (outside cooperative launch) one block at a time, so nothing can be reordered — every fence is already satisfied.	sim only	simulator.h:777
inline float __powf(float a, float b)		sim only	simulator.h:785	View
inline float __saturatef(float x)		sim only	simulator.h:796
inline void __sincosf(float x, float* s, float* c)		sim only	simulator.h:794
inline float __sinf(float x)		sim only	simulator.h:791
inline void __stcg(T* p, T v)		sim only	simulator.h:770
inline void __stcs(T* p, T v)		sim only	simulator.h:768	View
inline void __stwb(T* p, T v)		sim only	simulator.h:769
inline float __tanf(float x)		sim only	simulator.h:793
inline void __threadfence()		sim only	simulator.h:778	View
inline void __threadfence_block()		sim only	simulator.h:779
inline void __threadfence_system()		sim only	simulator.h:780
inline unsigned __umul24(unsigned a, unsigned b)		sim only	simulator.h:798	View
inline double rsqrt(double x)		sim only	simulator.h:784
inline float rsqrtf(float x)	Fast-math / integer intrinsics — computed exactly on the CPU (no reduced- precision device path), so the “__” fast variant equals the standard one.	sim only	simulator.h:783	View

Examples

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

__expf (test_web_cs_imagedenoising.cu:132)

        float  fCount     = 0;
        float  sumWeights = 0;
        float3 clr        = {0, 0, 0};
        float4 clr00      = tex2D<float4>(texImage, x, y);

        for (float i = -KNN_WINDOW_RADIUS; i <= KNN_WINDOW_RADIUS; i++)
            for (float j = -KNN_WINDOW_RADIUS; j <= KNN_WINDOW_RADIUS; j++) {
                float4 clrIJ      = tex2D<float4>(texImage, x + j, y + i);
                float  distanceIJ = vecLen(clr00, clrIJ);

                float weightIJ = __expf(-(distanceIJ * Noise
                                          + (i * i + j * j) * INV_KNN_WINDOW_AREA));

                clr.x += clrIJ.x * weightIJ;
                clr.y += clrIJ.y * weightIJ;
                clr.z += clrIJ.z * weightIJ;

                sumWeights += weightIJ;
                fCount += (weightIJ > KNN_WEIGHT_THRESHOLD) ? INV_KNN_WINDOW_AREA : 0;
            }

        sumWeights = 1.0f / sumWeights;

__ldcs (test_llmc_attention_backward.cu:479)

            local_sum += att_bth[t2] * datt_bth[t2];
        }

        block_acc[warp.meta_group_rank()] = cg::reduce(warp, local_sum, cg::plus<float>{});
        block.sync();
        local_sum = cg::reduce(warp, block_acc[warp.thread_rank()], cg::plus<float>{});

        for (int t3 = block.thread_rank(); t3 <= t; t3 += BlockSize) {
            // don't touch the cache. Some parts will still be here from the previous loop, and
            // we want to exploit those.
            float acc = __ldcs(att_bth + t3) * (__ldcs(datt_bth + t3) - local_sum);
            __stcs(dpreatt_bth + t3, scale * acc);
        }
    }
}

// ----------------------------------------------------------------------------
// kernel launchers ([SIM] <<<dim3(...),block>>> -> LAUNCH(kernel, dim3(...), block, ...))

void launch_softmax_1(float* dpreatt, float* datt, const float* att, int B, int T, int C, int NH, int block_size) {
    int num_blocks = ceil_div(T, block_size);
    LAUNCH(softmax_autoregressive_backward_kernel1, dim3(num_blocks, B*NH), block_size, dpreatt, datt, att, B, T, C, NH);

__ldg (test_web_transpose.cu:68)

        B[ny * N + nx] = A[nx * N + ny];
    }
}

__global__ void transpose3(const real *A, real *B, const int N)
{
    const int nx = blockIdx.x * blockDim.x + threadIdx.x;
    const int ny = blockIdx.y * blockDim.y + threadIdx.y;
    if (nx < N && ny < N)
    {
        B[ny * N + nx] = __ldg(&A[nx * N + ny]);
    }
}

// CPU reference: B is the transpose of the N x N matrix A.
static bool verify_transpose(const std::vector<real> &A, const std::vector<real> &B, int N)
{
    for (int i = 0; i < N; i++)
        for (int j = 0; j < N; j++)
            if (B[i * N + j] != A[j * N + i]) return false;
    return true;
}

__mul24 (test_web_cs_convtexture.cu:52)

    if (e != cudaSuccess) {
        std::printf("CUDA error in %s: %s\n", what, cudaGetErrorString(e));
        g_ok = false;
    }
}

#define KERNEL_RADIUS 8
#define KERNEL_LENGTH (2 * KERNEL_RADIUS + 1)

// ---- GPU-specific defines (verbatim) ----------------------------------------
#define IMAD(a, b, c) (__mul24((a), (b)) + (c))
#define UNROLL_INNER 1

static inline int iDivUp(int a, int b) { return (a % b != 0) ? (a / b + 1) : (a / b); }

// ---- Convolution kernel storage (verbatim) ----------------------------------
__constant__ float c_Kernel[KERNEL_LENGTH];

extern "C" void setConvolutionKernel(float* h_Kernel) {
    cudaMemcpyToSymbol(c_Kernel, h_Kernel, KERNEL_LENGTH * sizeof(float));
}

__powf (test_web_pcu_intrinsics.cu:70)

{
    int i;
    int tid = (blockDim.x * blockIdx.x) + threadIdx.x;

    if(tid == 0)
    {
        float tmp;

        for (i = 0; i < iters; i++)
        {
            tmp = __powf(a, 2.0f);
        }

        *out = tmp;
    }
}

int main(int argc, char **argv)
{
    (void)argc; (void)argv;

    // [SIM] graceful no-GPU guard (pattern: test_vector_add.cu)

__stcs (test_llmc_attention_backward.cu:480)

        }

        block_acc[warp.meta_group_rank()] = cg::reduce(warp, local_sum, cg::plus<float>{});
        block.sync();
        local_sum = cg::reduce(warp, block_acc[warp.thread_rank()], cg::plus<float>{});

        for (int t3 = block.thread_rank(); t3 <= t; t3 += BlockSize) {
            // don't touch the cache. Some parts will still be here from the previous loop, and
            // we want to exploit those.
            float acc = __ldcs(att_bth + t3) * (__ldcs(datt_bth + t3) - local_sum);
            __stcs(dpreatt_bth + t3, scale * acc);
        }
    }
}

// ----------------------------------------------------------------------------
// kernel launchers ([SIM] <<<dim3(...),block>>> -> LAUNCH(kernel, dim3(...), block, ...))

void launch_softmax_1(float* dpreatt, float* datt, const float* att, int B, int T, int C, int NH, int block_size) {
    int num_blocks = ceil_div(T, block_size);
    LAUNCH(softmax_autoregressive_backward_kernel1, dim3(num_blocks, B*NH), block_size, dpreatt, datt, att, B, T, C, NH);
}

__threadfence (test_web_cbe_hashtable.cu:51)

// [SIM] lock.h, inlined verbatim apart from HANDLE_ERROR staying local.
struct Lock {
    int* mutex;
    Lock() {
        HANDLE_ERROR(cudaMalloc((void**)&mutex, sizeof(int)));
        HANDLE_ERROR(cudaMemset(mutex, 0, sizeof(int)));
    }
    __device__ void lock() {
        while (atomicCAS(mutex, 0, 1) != 0);
        __threadfence();
    }
    __device__ void unlock() {
        __threadfence();
        atomicExch(mutex, 0);
    }
};

#define SIZE          (64 * 1024)                 // [SIM] was 100*1024*1024
#define ELEMENTS      (SIZE / sizeof(unsigned int))
#define HASH_ENTRIES  1024

__umul24 (test_web_cs_simpletexture3d.cu:57)

        g_ok = false;
    }
}

cudaArray*          d_volumeArray = 0;
cudaTextureObject_t tex;               // 3D texture

// ---- verbatim kernel from the sample ---------------------------------------
__global__ void d_render(uint* d_output, uint imageW, uint imageH, float w,
                         cudaTextureObject_t texObj) {
    uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
    uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;

    float u = x / (float)imageW;
    float v = y / (float)imageH;
    // read from 3D texture
    float voxel = tex3D<float>(texObj, u, v, w);

    if ((x < imageW) && (y < imageH)) {
        // write output color
        uint i      = __umul24(y, imageW) + x;
        d_output[i] = voxel * 255;

rsqrtf (test_web_pk_nbody.cu:44)

            __shared__ float3 shared_position[BLOCK_SIZE];
            float4 temp_position = p[tile * blockDim.x + threadIdx.x];
            shared_position[threadIdx.x] = make_float3(temp_position.x, temp_position.y, temp_position.z);
            __syncthreads();

            for (int j = 0; j < BLOCK_SIZE; j++) {
                float dx = shared_position[j].x - p[i].x;
                float dy = shared_position[j].y - p[i].y;
                float dz = shared_position[j].z - p[i].z;
                float distSqr = dx*dx + dy*dy + dz*dz + SOFTENING;
                float invDist = rsqrtf(distSqr);
                float invDist3 = invDist * invDist * invDist;

                Fx += dx * invDist3; Fy += dy * invDist3; Fz += dz * invDist3;
            }
            __syncthreads();
        }

        v[i].x += dt*Fx; v[i].y += dt*Fy; v[i].z += dt*Fz;
    }
}