PTX createpolicy instruction compile failure

Accelerated Computing CUDA CUDA NVCC Compiler
, ,
1

Hey, I’ve been trying to the use the createpolicy instruction from the examples in the PTX ISA documentation linked here (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html?highlight=createpolicy#data-movement-and-conversion-instructions-createpolicy). This is my code:

__device__ uint32_t ld_with_create_policy(const uint32_t * ptr) {
    uint32_t ret;
    asm volatile("createpolicy.fractional.L2::evict_last.b64 cache-policy, 0.5;");
    asm volatile("ld.global.nc.L2::cache_hint" ".u32 %0, [%1], cache-policy;" : "=r"(ret) : "l"(ptr));
    return ret;
}

__global__ void test_kernel(uint32_t * out, const uint32_t * in) {
    auto i = threadIdx.x + blockIdx.x * blockDim.x;
    uint32_t t = ld_with_create_policy(in + i);
    out[i] = t + 5;
}

This is the error I get:

/usr/local/cuda-12.9/bin/nvcc  -ccbin g++ -gencode=arch=compute_90,code=sm_90 -Xcompiler -O3  create_policy_test.cu -lcudart -o create_policy_test
ptxas /tmp/tmpxft_0014e837_00000000-6_create_policy_test.ptx, line 34; fatal   : Parsing error near '-': syntax error
ptxas fatal   : Ptx assembly aborted due to errors
make: *** [Makefile:10: create_policy_test] Error 255

NVCC version 12.9
GPU: H100

It seems to indicate that the - in cache-policy is an illegal syntax?

Can you let me know if I’m doing something wrong?

Thanks

2

cache-policy is the name of the operand, kind of like “destination” could be the “name” of an operand

none of the examples show usage of “cache-policy” as an actual operand

The createpolicy instruction creates a cache eviction policy for the specified cache level in an opaque 64-bit register specified by the destination operand cache-policy .

So we need a 64-bit register designation there.

3

Thanks! This works.